# onegpu docs, all pages (OneGpu 1.0.0) Every page of https://recursiveai.org/onegpu/docs/ as markdown, getting started first. Index: https://recursiveai.org/onegpu/llms.txt # Getting started with OneGpu Create a GPU device from C#, upload two matrices into typed buffers, multiply them with a compute kernel from `RecursiveAI.OneGpu.Kernels`, and read the result back. This guide was run on the 1.0.0 packages from NuGet, on Linux with Vulkan, on an NVIDIA GeForce RTX 4070. ## Install You need the .NET 10 SDK and, on Linux or Windows, a Vulkan 1.3 driver with the loader library (`libvulkan.so.1`, `vulkan-1.dll`); on macOS, Apple silicon. Make a console project and add the two packages: ```shell dotnet new console -n GpuHello cd GpuHello dotnet add package RecursiveAI.OneGpu.Backends --version 1.0.0 dotnet add package RecursiveAI.OneGpu.Kernels --version 1.0.0 ``` `RecursiveAI.OneGpu.Backends` brings the Vulkan, Metal and WebGPU backends and picks one for the platform: Metal on Apple platforms, WebGPU in a browser, Vulkan elsewhere. ## Multiply two matrices on the GPU Replace `Program.cs` with: ```csharp using RecursiveAI.OneGpu; using RecursiveAI.OneGpu.Kernels; // Every adapter the selected backend can see. Console.WriteLine($"Backend: {GpuBackends.Selected}"); foreach (var info in GpuBackends.EnumerateAdapters()) Console.WriteLine($" {info.Kind,-10} {info.Name}"); // Ask for the discrete GPU by kind; a name fragment such as "RTX" works too. using var device = GpuBackends.CreateDevice( new GpuDeviceOptions { Adapter = "discrete" }); var adapter = device.Adapter; Console.WriteLine( $"Running on: {adapter.Name} ({adapter.Kind}, API {adapter.ApiVersion})"); // C = A x B: A is 4 x 2, B is 2 x 3, both row-major. const int M = 4, N = 3, K = 2; float[] a = [1, 2, 3, 4, 5, 6, 7, 8]; float[] b = [1, 0, 2, 0, 1, 3]; // Typed buffers; A and B are uploaded before the first submission. using var bufferA = device.CreateBuffer(a); using var bufferB = device.CreateBuffer(b); using var bufferC = device.CreateBuffer(M * N); using var gemm = new Gemm(device, GemmElementType.F32); var encoder = device.BeginCommands(); gemm.Encode(encoder, new GemmProblem(M, N, K), bufferA, bufferB, bufferC); device.Submit(encoder); // Completes when the GPU is done, without blocking a thread. var c = new float[M * N]; await bufferC.ReadAsync(c); for (int row = 0; row < M; row++) Console.WriteLine(string.Join(" ", c.Skip(row * N).Take(N).Select(v => v.ToString().PadLeft(3)))); ``` Run it: ```shell dotnet run ``` The output, on a Linux machine with an RTX 4070 and a Radeon integrated GPU: ```text Backend: Vulkan Discrete NVIDIA GeForce RTX 4070 Integrated AMD Ryzen 9 7950X3D 16-Core Processor (RADV RAPHAEL_MENDOCINO) Software llvmpipe (LLVM 21.1.8, 256 bits) Running on: NVIDIA GeForce RTX 4070 (Discrete, API 1.4.312) 1 2 8 3 4 18 5 6 28 7 8 38 ``` The first row is `[1 2]` times B: 1·1 + 2·0 = 1, 1·0 + 2·1 = 2, 1·2 + 2·3 = 8. ## What each part does - **`GpuBackends.CreateDevice`** opens a device on the backend for this platform. `GpuDeviceOptions.Adapter` picks the adapter whose name contains the text, or whose kind is `software`, `discrete` or `integrated`; without it the backend takes the first adapter in its order, which on Vulkan lists discrete GPUs first, then integrated, virtual and software ones. The `ONEGPU_ADAPTER` environment variable does the same and wins over the option. `device.Adapter` reports the adapter the device opened. - **`CreateBuffer(a)`** makes a typed `GpuBuffer` and queues the upload of `a`, which runs before the next submission. `CreateBuffer(M * N)` makes an empty one of that many elements. - **`Gemm`** is a kernel family: create it once and it keeps its pipelines. `Encode` records the compute passes for one product into the encoder, and `Submit` sends them to the GPU. Pass `GpuKernelMode.Deterministic` as its third argument for bit-repeatable sums on one GPU and driver; the default is `Fast`. - **`ReadAsync`** copies the result back without blocking a thread; it completes when the GPU has finished the work submitted before it. - A wrong shape, stride, size or offset throws `GpuKernelShapeException` before anything reaches the driver; a missing capability throws `GpuUnsupportedException`. ## Next - [Compute surface](https://recursiveai.org/onegpu/docs/compute/): capabilities, buffers and readback, timestamps, passes and ordering - [Kernels](https://recursiveai.org/onegpu/docs/kernels/): every operation, with its tolerances - [Matrix multiplication (GEMM)](https://recursiveai.org/onegpu/docs/gemm/) - [Names and environment variables](https://recursiveai.org/onegpu/docs/environment/): `ONEGPU_BACKEND`, `ONEGPU_ADAPTER` and validation # onegpu GPU compute and rendering for .NET from pure C#, on Vulkan, Metal and browser WebGPU. There are no native binaries of ours and no CUDA install: the backends talk to the platform's own GPU API through the driver already on the machine (the Vulkan loader, the Objective-C runtime and Metal framework, the browser's WebGPU). It has two jobs. It is a small, compute-first device layer (typed buffers, async readback, compute pipelines, command encoders, timestamps, capabilities as data) with the surfaces, textures and graphics pipelines a renderer needs. And it ships the compute kernels that machine-learning inference needs, written in HLSL and compiled at build time for all three backends. onegame, the family's game engine, takes its GPU layer from these packages (its render graph and 2D/3D rendering stay in onegame). OneModel's GPU support, which runs ONNX models on these kernels, is coming in a later OneModel release. **Status:** 1.0.0, released on [nuget.org](https://www.nuget.org/packages/RecursiveAI.OneGpu.Backends/1.0.0). The public API is frozen: a change that breaks it needs a new major version. | Package | What it is | |---|---| | `RecursiveAI.OneGpu` | Devices, adapters and capabilities, typed buffers, uploads and async readback, compute pipelines, command encoders, dispatch, timestamps, memory statistics, plus surfaces, textures and graphics pipelines. BCL only; no interop. | | `RecursiveAI.OneGpu.Vulkan` | Vulkan 1.3 backend (Windows, Linux). | | `RecursiveAI.OneGpu.Metal` | Metal backend (macOS, Apple silicon). | | `RecursiveAI.OneGpu.WebGpu` | WebGPU backend for the browser (`[JSImport]`, publishes `onegpu.js`). | | `RecursiveAI.OneGpu.Backends` | `GpuBackends`: picks Metal on Apple platforms, WebGPU in a browser and Vulkan elsewhere. Reference it to get all three. | | `RecursiveAI.OneGpu.Kernels` | HLSL compute kernels with generated bindings: GEMM, attention, convolution, reductions, softmax, norms, gather and scatter, int8, recurrent layers, element-wise, integer and bool ops. | ## Platforms and requirements - **.NET 10.** Every package is trimmable and Native AOT compatible. - **Linux and Windows:** a Vulkan 1.3 driver (NVIDIA, AMD, Intel; the Mesa software driver works too). Tested on NVIDIA and AMD hardware and on llvmpipe; Intel is not tested. The loader library (`libvulkan.so.1`, `vulkan-1.dll`) must be installed. - **macOS:** Apple silicon with Metal. Tested on an M2. - **Browser:** Chrome with WebGPU. `shader-f16` and `subgroups` are used when the adapter offers them; a kernel that needs a missing feature falls back to a portable variant or reports it by name. - No OpenGL, no CUDA, no mobile. ## Install ``` dotnet add package RecursiveAI.OneGpu.Backends --version 1.0.0 dotnet add package RecursiveAI.OneGpu.Kernels --version 1.0.0 ``` `Backends` brings the three backends; reference `RecursiveAI.OneGpu` and one backend alone to be smaller. ## A compute example ```csharp using RecursiveAI.OneGpu; using RecursiveAI.OneGpu.Kernels; using var device = GpuBackends.CreateDevice(); // Metal on Apple silicon, Vulkan elsewhere Console.WriteLine($"{device.Adapter.Name}: {device.Capabilities.Compute.Describe()}"); const int M = 4, N = 3, K = 2; float[] a = [1, 2, 3, 4, 5, 6, 7, 8]; // 4 x 2, row-major float[] b = [1, 0, 2, 0, 1, 3]; // 2 x 3, row-major using var bufferA = device.CreateBuffer(a); using var bufferB = device.CreateBuffer(b); using var bufferC = device.CreateBuffer(M * N); using var gemm = new Gemm(device, GemmElementType.F32); // GpuKernelMode.Deterministic for bit-repeatable sums var encoder = device.BeginCommands(); gemm.Encode(encoder, new GemmProblem(M, N, K), bufferA, bufferB, bufferC); device.Submit(encoder); var c = new float[M * N]; await bufferC.ReadAsync(c); // completes when the GPU is done, without blocking a thread ``` A family such as `Gemm` is created once and keeps its pipelines; `Encode` records one or more compute passes into the encoder you pass and allocates nothing once its bind groups are cached. Your own shaders work the same way: compile an HLSL file with `RecursiveAI.OneShader.Build` and pass its generated entry point to `GpuDevice.CreateComputeKernel`. [`docs/compute.md`](https://recursiveai.org/onegpu/docs/compute/) has the details: capabilities, timestamps, indirect dispatch, memory statistics. Bad input is checked in safe code before anything reaches the driver: a wrong shape, stride, size or offset throws `GpuKernelShapeException` naming the operation and the argument, and a missing capability or exceeded limit throws `GpuUnsupportedException` with `Capability`, `Expected` and `Actual`. ## Kernel families | Family | Operations | |---|---| | `Gemm`, `Int8Gemm` | batched matrix products in float32, float16 and bfloat16 (cooperative-matrix variants where the device has them), uint8/int8 with int32 sums | | `AttentionKernels` | fused scaled dot-product attention with masks, causal rule and grouped heads | | `ConvKernels`, `ConvIntegerKernels`, `PoolKernels`, `ResizeKernels` | Conv, ConvTranspose, ConvInteger, pooling, resize | | `RecurrentKernels` | RNN, GRU, LSTM | | `ElementwiseKernels`, `IntegerKernels`, `ReduceKernels`, `SoftmaxKernels`, `NormKernels` | broadcasting binary and unary ops, integer and bool ops, reductions, softmax, layer and RMS normalisation | | `GatherKernels`, `ScatterKernels`, `LayoutKernels`, `TransposeKernels` | gather and scatter, slice, concat, expand, transpose | | `DequantizeKernels`, `DynamicQuantizeKernels` | QuantizeLinear, DequantizeLinear, DynamicQuantizeLinear | [`docs/kernels.md`](https://recursiveai.org/onegpu/docs/kernels/) lists every operation; the other pages in [`docs/`](https://github.com/recursive-ai-oss/onegpu-csharp/tree/5d3317caecb82fedd96642e54b0ab33bcb9759c5/docs) cover GEMM, attention, convolution and recurrent layers in detail. ## Determinism and tolerances - **Fast is the default.** A fast kernel may sum in any order and use fused multiply-add. - **`GpuKernelMode.Deterministic`** exists for GEMM, reductions, softmax and attention: a fixed reduction and accumulation order, no float atomics, nothing that depends on occupancy or on how the driver schedules subgroups. Two runs on the same GPU and driver give identical bits. Results are not bit-identical across vendors or drivers. Other families give the same bits in both modes. - **Accuracy.** The tests compare every family with a plain C# reference in double precision. Integer, bool, copy and cast operations, float32 add, sub, mul, min and max, and int8 GEMM are bit-exact. Float32 GEMM, reductions, softmax, norms and transcendentals are within 1e-5 of the reference, scaled by the largest output value (measured values are far inside that); float16 inputs with float32 accumulation within 2e-3; bfloat16 within 8e-3. The limit per operation and dtype is in [`docs/kernels.md`](https://recursiveai.org/onegpu/docs/kernels/). - **Capabilities are queried, never assumed.** Each capability-gated variant has a portable fallback that the tests also run. ## Browser Reference `RecursiveAI.OneGpu.Backends` (or `.WebGpu`) from a .NET WebAssembly browser app. The package publishes `onegpu.js` beside the app. Create the device with `await GpuBackends.CreateDeviceAsync()` and read results with `ReadAsync`, because the browser cannot block. The browser records each submission as one byte stream and runs it in a single call into the script. Float16 and bfloat16 data work without the `shader-f16` feature (the kernels read them as packed 32-bit words). The kernel families' bind-group caches (`maxCachedBindGroups`) must cover everything one command buffer binds, because the browser keeps those groups alive until the submission runs. Fused attention has browser-specific variants; see [`docs/attention.md`](https://recursiveai.org/onegpu/docs/attention/). Environment variables, the JavaScript module's surface and how the public API is frozen are in [`docs/environment.md`](https://recursiveai.org/onegpu/docs/environment/). ## Speed > **Note:** these figures were “Measured through onemodel’s `RecursiveAI.OneModel.Gpu`” (below). That package isn’t released yet; OneModel’s GPU support is coming in a later OneModel release, so the figures can’t be reproduced with the released packages. Measured through onemodel's `RecursiveAI.OneModel.Gpu` on onemodel's reference set (MiniLM-L6, BERT-base, ResNet-50, MobileNetV3, DBNet, CRNN, Whisper-small), against ONNX Runtime's GPU providers on the same machine and model file. Ratio is ours divided by the reference; the bar is 1.5 (target 1.15) on a desktop and 2.5 (target 1.5) in a browser. Full tables, method and the correctness runs: [`docs/speed.md`](https://recursiveai.org/onegpu/docs/speed/). | Platform | Reference | Rows judged | Result | |---|---|---|---| | Linux, RTX 4070, Vulkan | ONNX Runtime 1.31.1, WebGPU provider | 14 | all within the bar; worst ResNet-50 b8 1.40, BERT b8 1.35; the rest at or under 1.15 | | macOS, Apple M2, Metal | whichever of Core ML (GPU) and the WebGPU provider took less time | 14 | 12 within the bar; **ResNet-50 float32 misses: 2.16 at batch 1 (13.09 vs 6.07 ms) and 2.21 at batch 8 (77.94 vs 35.27 ms) against Core ML on the GPU**, 0.82 and 0.81 against the WebGPU provider on Metal | | Chrome 155, RTX 4070, WebGPU | onnxruntime-web 1.31.1 | 12 | all within the bar; 11 within the target | | Windows 11, Intel Iris Xe, Vulkan (indicative) | ONNX Runtime 1.24.4, DirectML | 20 | misses on float32 BERT, MiniLM, ResNet-50 and Whisper, float16 ResNet-50 and MobileNetV3 at batch 8 (1.53 to 2.89; this GPU has no cooperative matrix); MobileNetV3 b1, DBNet, CRNN, float16 MiniLM and every 8-bit row within the bar | The misses are known and ship as they are; rows where the reference leaves compute to its CPU (the 8-bit models) are listed for information only. Native AOT on Windows has not been tested. ## Build and test ``` dotnet build OneGpu.slnx tools/ci/test-vulkan.sh # Vulkan, validation on tools/ci/test-metal.sh # macOS tools/ci/test-web.sh # headless Chrome on WebGPU tools/ci/test-web-package.sh # a browser app consuming the packed WebGPU backend, in headless Chrome tools/ci/aot-smoke.sh # Native AOT publish of a small program, run on a device ``` Output goes to `artifacts/` at the repository root, or to `$ONEGPU_ARTIFACTS` when set. Device tests skip, and print why, when no adapter is present. On Vulkan, run with the Khronos validation layer (`VK_INSTANCE_LAYERS=VK_LAYER_KHRONOS_validation`, or `ONEGPU_VALIDATION=1`); a run that prints validation messages fails. On Metal the scripts turn on API and shader validation. ## Licence `MIT OR Apache-2.0`, © Recursive AI, LLC. See [`THIRD-PARTY-NOTICES.md`](https://github.com/recursive-ai-oss/onegpu-csharp/blob/5d3317caecb82fedd96642e54b0ab33bcb9759c5/THIRD-PARTY-NOTICES.md). # Compute surface What `RecursiveAI.OneGpu` offers for compute, how each backend fills it in, and the rules that hold on every backend. ## Capabilities `GpuDevice.Capabilities.Compute` is a `GpuComputeCapabilities`. Every value comes from the graphics API's own query on the adapter, never from the vendor name or the API version. Zero means the API did not report the value. | Value | Vulkan | Metal | WebGPU | |---|---|---|---| | Subgroup size, min, max | `VkPhysicalDeviceSubgroupProperties`, `SubgroupSizeControlProperties` | `threadExecutionWidth` of a compiled probe pipeline | `GPUAdapterInfo.subgroupMinSize/MaxSize` with the `subgroups` feature; size itself is not reported | | Subgroup operations | `supportedOperations`, when the compute stage is listed | one probe kernel per operation class must compile | all but clustered, when the `subgroups` feature is on | | Shader features (`DeviceFeatures`) | core 1.1/1.2 feature structs, `shaderInt16/64`, `shaderFloat64` | half, 16-bit and 64-bit integers are available on every Apple GPU | `shader-f16` | | 8-bit integers | `shaderInt8` and `storageBuffer8BitAccess` | yes | no | | Packed int8 dot product | `shaderIntegerDotProduct` | no (Metal has no packed dot instruction) | WGSL language feature `packed_4x8_integer_dot_product` | | bfloat16 | `VK_KHR_shader_bfloat16` | a `bfloat` kernel compiles (Metal Shading Language 3.1) | no | | Cooperative matrices | `VK_KHR_cooperative_matrix` properties, subgroup scope | `simdgroup_matrix` 8×8×8 combinations whose probe kernel compiles | none | | Full subgroups (`FullSubgroups`) | Vulkan 1.3 `computeFullSubgroups` | yes (a threadgroup that is a multiple of the SIMD-group width runs whole SIMD groups) | no | | Selectable subgroup size (`SubgroupSizeControl`) | Vulkan 1.3 `subgroupSizeControl` and the compute stage in `requiredSubgroupSizeStages` | no (fixed at 32 on Apple GPUs) | no | | Workgroup size, invocations, count, memory | `VkPhysicalDeviceLimits` | `maxThreadsPerThreadgroup`, `maxTotalThreadsPerThreadgroup` of the probe, `maxThreadgroupMemoryLength` | `maxCompute*` limits of the device | | Buffer limits | `maxStorageBufferRange`, `maxMemoryAllocationSize` | `maxBufferLength` | `maxStorageBufferBindingSize`, `maxBufferSize` | Metal has no feature list for its shader language, so a feature counts when a kernel that uses it compiles on the device. The probes run once, the first time `Capabilities` is read. Vulkan devices enable the features they report when they are created (16-bit and 8-bit storage, `shaderFloat16`, `shaderInt8`, integer dot product, subgroup size control, cooperative matrix, bfloat16), so a kernel that was checked against the capabilities can use them. A WebGPU device requests `shader-f16` and `subgroups` when the adapter has them. ## Kernels `GpuDevice.CreateComputeKernel(in ShaderSource, label)` builds the shader, the pipeline layout and the pipeline from the entry point's reflection. Before it creates anything it checks the kernel against the device and throws `GpuUnsupportedException` naming the missing capability (`Capability`, `Expected`, `Actual`): the workgroup size against the per-axis and total limits, the required scalar types against `ShaderFeatures`, the subgroup operations against `SubgroupOperations`. The bindings and push-constant block come from the reflection, so the host cannot declare a layout that differs from the shader's. **Workgroup memory a pipeline uses.** `GpuComputePipeline.ReportedWorkgroupMemoryBytes` is what the compiled pipeline reports (Metal's `staticThreadgroupMemoryLength`; 0 on backends that do not report it). Metal enforces its limit against that figure, which is larger than the declared size while shader validation is on (it doubles it) and when a three-component vector takes 16 bytes, so a kernel with several shapes picks one from this number after compiling (the float32 GEMM does). **Subgroup requirements.** `ComputePipelineDesc.RequireFullSubgroups` makes every subgroup of every workgroup full (the workgroup size must be a multiple of the subgroup size) and `RequiredSubgroupSize` pins the width. Cooperative-matrix kernels need the first below SPIR-V 1.6, so the shader compiler lists `FullSubgroups` in the reflection of every entry point that uses them and `CreateComputeKernel` sets it for any kernel whose reflection lists that feature (devices report it in `ShaderFeatures` when `GpuComputeCapabilities.FullSubgroups` is true), and takes an optional `requiredSubgroupSize` for kernels written for one width. Vulkan passes `VK_PIPELINE_SHADER_STAGE_CREATE_REQUIRE_FULL_SUBGROUPS_BIT` and chains `VkPipelineShaderStageRequiredSubgroupSizeCreateInfo` (Vulkan 1.3 core, `VK_EXT_subgroup_size_control`). A request the device cannot meet throws `GpuUnsupportedException` whose `Capability` is `FullSubgroups` or `RequiredSubgroupSize`, before anything is created, on every backend: a device with one fixed subgroup size accepts that size and no other; Metal accepts full subgroups (its SIMD groups are whole when the threadgroup is a multiple of the SIMD width, which `CreateComputeKernel` checks against the workgroup size of a cooperative-matrix kernel); WebGPU has no control over either and rejects full subgroups and any pinned size unless the adapter reports a single size. Push constants work on every backend. On WebGPU they are a uniform buffer in the shader's highest bind group, written into a ring and bound with a dynamic offset; a bind group that the application makes for that group carries the binding. ## Buffers and readback `GpuBuffer` wraps a `GpuBuffer` of `T` elements (`T` unmanaged; its layout must match the shader's). `Write` queues an upload that runs before the next submission. `Read` waits for the device. `ReadAsync` (and `GpuDevice.ReadBufferAsync` for bytes) records a copy into a host-visible buffer behind the work submitted so far, and a helper thread waits for that submission alone; the calling thread and the device keep working, later submissions are not delayed, and the returned `ValueTask` completes when the bytes are in the destination. On WebGPU the browser's own asynchronous map does the same. Completions arrive in submission order. Disposing the device finishes every pending readback first. ## Timestamps `Capabilities.Timestamps` says whether passes can be timed. `BeginPassTiming` names a `GpuTimestampSet` and the next passes write their start and end into consecutive slots; `TryReadTimestamps` returns nanoseconds once the submission has finished. - **Vulkan** writes a timestamp query at each end of the pass; the resolution is the device's `timestampPeriod`. - **Metal** samples the GPU clock at the boundaries of a compute pass or render pass (counter sampling at stage boundaries, which Apple GPUs support; they cannot sample between draws or dispatches, so one pass is the finest unit). The samples are ticks of the GPU clock; the device converts them to nanoseconds with the rate it measures between the CPU and GPU clocks. On an Apple M2 one tick is 1 ns, and the compute pass times of a test run add up to within 2 % of the wall time of the run that held them. Metal skips the sample of a stage that did no work, so a render pass without draws or a compute pass without dispatches reads as zero length at the time of the pass before it. - **WebGPU** does not report timestamps. ## Passes and ordering Compute passes recorded one after another in one `GpuCommandEncoder` run in that order: a pass sees every write of the passes before it, and a later write to a buffer wins. Copies recorded between passes are ordered the same way. - **Vulkan** records every pass into the same command buffer and separates passes with memory barriers; a pass costs nothing beyond its commands. - **Metal** keeps one compute encoder open across consecutive passes. An encoder created by `computeCommandEncoder` is serial (`MTLDispatchTypeSerial`): each dispatch runs to completion, with its writes visible, before the next starts, which is the ordering above. A copy, a render pass or the end of the recording closes the encoder. A pass that is timed gets an encoder of its own, because counter samples are taken at encoder boundaries; timing a chain of tiny passes therefore costs about 20 µs per pass on an Apple M2 (the sample buffers), where the same chain untimed costs 3 µs per dependent dispatch. Time whole recordings, or the few passes you care about, rather than every pass of a long chain. - **WebGPU** ends and begins a compute pass in the browser's own encoder; the browser inserts the barriers. `DispatchChainTests` records chains of dependent passes (transposes of a matrix back and forth, repeated full writes to one buffer, copies between passes) and checks every element. ### Counting passes `GpuCommandEncoder.ComputePassCount` and `DispatchCount` say how many compute passes were begun and how many `Dispatch` and `DispatchIndirect` calls were recorded since `BeginCommands` handed the encoder out (read them before `Submit`). The encoder counts as it records, so they need no timestamps and cost nothing, and they count what the caller recorded on every backend: Metal's shared compute encoders do not change them. An empty grid (a zero dimension) is counted although it runs nothing. `PassCountTests` checks the counts on every backend. ## Errors | Exception | Raised when | Extra data | |---|---|---| | `GpuUnsupportedException` | a capability is missing or a limit is exceeded | `Capability`, `Expected`, `Actual` | | `GpuOutOfMemoryException` | the device or host has no memory left | `RequestedBytes` | | `GpuDeviceLostException` | the driver reset or the GPU was removed | `Operation`, the API call that saw it | ## Memory statistics `GpuDevice.MemoryStats` covers compute buffers like any other buffer. Vulkan reports the blocks held from the driver, the bytes live resources occupy and the staging memory; resources that were disposed count until the GPU has finished with them. Metal reports the device's allocated size as reserved and the live buffers (count and bytes) as used; textures are not itemised. WebGPU keeps no tally and reports zeros. ## Devices and instances The Vulkan backend shares one instance per process (one more when validation is requested) and keeps it for the life of the process; each device is its own logical device on it. Some drivers refuse new instances after a few dozen create and destroy pairs in one process (NVIDIA 595 stopped after about 80), and a long-running tool or test run must not hit that. ## Adapter order The Vulkan backend keeps adapters in one total order, whatever order the drivers are enumerated in: kind (discrete, integrated, virtual, software, other), then PCI vendor id, device id, name, API version (newest first) and the driver's pipeline-cache UUID. The default is the first adapter in that order, and `GpuDeviceOptions.Adapter` or `ONEGPU_ADAPTER` (a name substring or a kind) picks the first match in it, so a multi-GPU machine chooses the same adapter every run. ## Validation messages `GpuDevice.TakeValidationMessages` returns what the layer said about this device. Devices share one instance and one message callback, so the Vulkan backend records which device created each object and sends a message to the owner of the first object it names; messages raised while a `VkDevice` is created or destroyed go to that device. Two devices alive at once never see each other's messages. A message that names no object any device owns goes to every live device, because dropping it would hide a real problem. Call `TakeValidationMessages` again after `Dispose`: destroying a Vulkan device is where the layer lists the objects that were never disposed (`vkDestroyDevice(): ... has N leaked objects`), and the Metal backend adds a line when buffers were still alive. The test fixtures do this, so every suite also checks for leaks. A command encoder that was begun and never submitted does not leak: the device destroys every encoder it made. ## Browser `RecursiveAI.OneGpu.WebGpu` ships `onegpu.js` as a static web asset. An app that references the package publishes it at `_content/RecursiveAI.OneGpu.WebGpu/onegpu.js`, and `WebGpu.CreateDeviceAsync` loads it and registers it as the module `onegpu`; nothing is copied by hand. A page that registers its own copy (`setModuleImports('onegpu', ...)`) after importing it keeps working: the script marks the page, and the library then does not import a second one. `tools/ci/test-web-package.sh` checks the whole path with `tests/RecursiveAI.OneGpu.WebConsumer`, an app that restores the packed libraries from a scratch feed. Several devices can share a page: each `CreateDeviceAsync` makes a context of its own in `onegpu.js` (its WebGPU device, handle table, push-constant ring and pass state), and every call from a device names its context. Buffers and bind groups belong to the device that made them, as in WebGPU. Disposing a bind group, buffer-backed or not, while a command encoder that recorded it has not been submitted is safe on every backend. On the browser backend an encoder pins each bind group it records and a group disposed meanwhile is destroyed in the stream right after that encoder's commands, because WebGPU forbids destroying an object queued commands still name. An encoder that is never submitted keeps its groups alive. ## Known environment issue With the NVIDIA Vulkan driver of the 595 series on Linux, loading the validation layer first leaves too little static thread-local storage for the driver, and instance creation fails with "cannot allocate memory in static TLS block". Setting `GLIBC_TUNABLES=glibc.rtld.optional_static_tls=4096` for the process fixes it; `tools/ci/env.sh` does. With the RADV driver on a Ryzen 7000 integrated GPU (Mesa 26.0.3, kernel 7.0.0-42), a dispatch that starts after the GPU has idled for a few milliseconds sometimes writes one wrong value: about 3% of dispatches of a 1536-workgroup shared-memory reduction, none while another process keeps the GPU at its high clock. A plain Vulkan program (no library code) reproduces it with every barrier in place, and the same GPU model passes on an older Mesa, kernel and firmware, so it is not a synchronisation fault in this library. `tools/gpu-idle-repro` holds the program and the numbers. Results from that GPU on that stack are not trusted until it is updated. # Names that are part of the interface Everything on this page is frozen for 1.x: a rename is a breaking change. ## Environment variables Read by the packages: | Variable | Read by | Meaning | |---|---|---| | `ONEGPU_BACKEND` | `GpuBackends` | `vulkan` or `metal` (any case) overrides the platform's choice. Anything else, or unset: Metal on macOS and iOS, Vulkan elsewhere. In a browser the backend is always WebGPU. | | `ONEGPU_ADAPTER` | Vulkan, Metal | Picks the adapter whose name contains this text (case-insensitive) or whose kind is `software`, `discrete` or `integrated`. Wins over `GpuDeviceOptions.Adapter`. The first match in the backend's adapter order is used; no match throws `GpuUnsupportedException`. | | `ONEGPU_VALIDATION` | Vulkan | `1` or `true` turns on the Khronos validation layer, as `GpuDeviceOptions.Validation` does. Messages are collected per device (`GpuDevice.TakeValidationMessages`). | | `ONEGPU_VALIDATION_OPTIONAL` | Vulkan | `1` or `true` makes `GpuDeviceOptions.RequireValidation` fall back to running without the layer, with a notice, when the layer is not installed. Otherwise a missing layer throws. | | `ONEGPU_VULKAN_LIBRARY` | Vulkan | A path or name of a Vulkan loader to try before the system's (`libvulkan.so.1`, `vulkan-1.dll`, `libvulkan.1.dylib`). | Read by the tests, tools and benchmarks only (never by a package): | Variable | Meaning | |---|---| | `ONEGPU_ARTIFACTS` | Build output folder (default `artifacts/` in the repository). Read by `Directory.Build.props`. | | `ONEGPU_REQUIRE_DEVICE` | `1` makes a missing GPU fail a device test instead of skipping it. `tools/ci/env.sh` sets it. | | `ONEGPU_CHROME` | The Chrome executable `tools/web/run-page.mjs` starts. | | `ONEGPU_PARITY_REPORT`, `ONEGPU_REPORT`, `ONEGPU_GOLDEN_REPORT` | A file that the parity, throughput and golden-image tests append their result rows to. | | `ONEGPU_UPDATE_GOLDENS`, `ONEGPU_CAPTURE_DIR` | `1` rewrites the golden images in `testdata/`; the folder the render tests save captures to. | | `ONEGPU_GEMM_FULL`, `ONEGPU_CONV_FULL`, `ONEGPU_ATTN_LARGE` | `1` adds the large shape sets, which take long on the CPU reference. | | `ONEGPU_GEMM_NO_HARDWARE`, `ONEGPU_GEMM_SUBGROUPS`, `ONEGPU_CONV_BENCH_NOCOOP`, `ONEGPU_CONV_BENCH_NOSPLIT` | Force a variant off, to run the fallback on a device that has the feature, or to compare. | | `ONEGPU_BENCH_VALIDATION` | `0` times without the validation layer. | | `ONEGPU_GEMM_BENCH`, `ONEGPU_CONV_BENCH`, `ONEGPU_CONV_BENCH_BATCHES`, `ONEGPU_CONV_BENCH_LAYERS`, `ONEGPU_CONV_BENCH_MODE`, `ONEGPU_CONV_BENCH_ONLY`, `ONEGPU_ATTN_BENCH`, `ONEGPU_RECURRENT_BENCH`, `ONEGPU_DISPATCH_BENCH` | Select and shape the throughput tests (they are skipped unless set). | The Vulkan loader's and drivers' own variables also apply: `VK_ICD_FILENAMES` picks a driver, `VK_INSTANCE_LAYERS` and `VK_ADD_LAYER_PATH` control layers. `tools/ci/fetch-validation-layer.sh [folder]` downloads the validation layer the hosted CI runs (Ubuntu 24.04's `vulkan-validationlayers`, version pinned in that script and read from it by `ci.yml`) into a folder; point `VK_ADD_LAYER_PATH` at it to reproduce CI locally. A second run against the newest Khronos release catches what the older layer does not model. On Metal, `MTL_DEBUG_LAYER=1`, `MTL_DEBUG_LAYER_ERROR_MODE=assert` and `MTL_SHADER_VALIDATION=1` make misuse fatal. ## The JavaScript module `onegpu` The WebGPU backend imports a script as the module named `onegpu` (`[JSImport("...", "onegpu")]`). The package publishes it at `_content/RecursiveAI.OneGpu.WebGpu/onegpu.js`; `WebGpu.CreateDeviceAsync` imports it from there unless the page already imported it and registered it under that name, which it signals by setting `globalThis.__onegpuLoaded`. The script's exports are an implementation interface between the package's managed code and the browser, not something an app calls. They are named here because a page that registers its own copy (for example to bundle it) must provide them: | Export | Purpose | |---|---| | `init(requestJson)` | Asynchronously requests an adapter and device and creates a context. Returns JSON with the context number, adapter name, limits, features and subgroup sizes. `{"fallback":true}` asks for the software adapter. | | `exec(ctx, bytes)` | Runs one recorded stream of resource creation, uploads and commands in the context. | | `takeMessages(ctx)` | Returns and clears the validation and uncaptured errors the context saw. | | `whenIdle(ctx)` | Resolves when the device's queue has finished all submitted work. | | `readBuffer(ctx, id, offset, size)`, `readTexture(ctx, id, mip, layer, width, height, depth, bytesPerTexel)` | Start a readback and resolve to a ticket. | | `takeRead(ticket, destination)` | Copies the bytes of a finished readback out. | | `configureCanvas(ctx, canvasId, width, height)`, `setCanvasSize(canvasId, width, height)` | Attach and resize a canvas for a surface. | Every device in a page is a context of its own, so several devices can share one page; each function takes the context number first (the canvas functions after the first take it implicitly through the canvas id). ## The public API baseline Each package has `PublicAPI.Shipped.txt` beside its project: every public type and member, written as the compiler sees it. The `Microsoft.CodeAnalysis.PublicApiAnalyzers` analyzer (build time only, `PrivateAssets="all"`) makes a build fail when the code and the file disagree, so a public change cannot happen without a change to the file in the same diff. To change the API, edit the code and add or remove the matching lines; the analyzer's messages (`RS0016`, `RS0017`) give the exact text of each missing line, and a code fix in an IDE writes them. A removal or a changed signature is a breaking change and needs a major version. `RecursiveAI.OneGpu.Kernels` also contains one generated class per HLSL file (namespace `RecursiveAI.OneGpu.Kernels.Shaders`, produced by the shader build), which is why its file is large: a changed kernel entry point or variant axis changes that surface and shows in the diff. Those classes are the shader blobs and layouts the hand-written families use; apps should use the families. # Kernels `RecursiveAI.OneGpu.Kernels` holds the compute kernels that machine-learning inference needs, written in HLSL and compiled by the shader build for SPIR-V (Vulkan), Metal Shading Language and WGSL (browser). This page describes the element-wise, reduction, softmax, normalisation and layout families; each is a public class created from a `GpuDevice`. ## Using a family ```csharp using var elementwise = new ElementwiseKernels(device); var encoder = device.BeginCommands(); elementwise.EncodeBinary(encoder, GpuBinaryOp.Add, x, 0, [rows, cols], bias, 0, [cols], y, 0, new GpuActivation(GpuUnaryOp.GeluErf)); device.Submit(encoder); ``` - A family is created once; pipelines are built the first time a variant is used and kept. Calling `Encode…` records one compute pass into the caller's encoder. With the same buffers, a call allocates nothing (bind groups are cached by buffer identity; offsets and sizes travel in push constants). - Operands are a `GpuBuffer`, an element offset and a shape. Tensors are dense and row-major. `T` is `float` or `Half` for arithmetic; `int`, `uint`, `long` where an operation only moves or converts bits. Comparison results and conditions are bool tensors (see below). - Every size, shape, axis, permutation and offset is checked in safe code before anything is recorded. A bad argument throws `GpuKernelShapeException`, which names the operation and the argument or dimension (`Reduce: axes: axis 3 is outside [-3, 3)`). - Families are not thread-safe. Disposing a family releases its pipelines once the GPU has finished with them. - `GpuKernelMode.Fast` is the default. `GpuKernelMode.Deterministic` is accepted by every family; for the element-wise and layout families there is no reduction order, so both modes give the same bits. ## Families | Family | Operations | |---|---| | `ElementwiseKernels` | binary Add, Sub, Mul, Div, Pow, Min, Max with ONNX broadcasting (up to 6 dimensions after merging, inputs up to rank 8) and an optional fused unary activation; unary Identity, Neg, Abs, Sqrt, Reciprocal, Exp, Log, Tanh, Erf, Sigmoid, Relu, Gelu (erf and tanh forms), Silu, HardSigmoid, HardSwish, Clip, LeakyRelu; Equal, Less, Greater, LessOrEqual, GreaterOrEqual over float and Half; Where over float, Half, int, uint and long; Cast between float, Half, int, long and bool | | `ReduceKernels` | Sum, Mean, Max, Min, Prod, L1, L2, SumSquare, LogSum over any set of axes; ArgMax and ArgMin (64-bit result) | | `SoftmaxKernels` | Softmax and LogSoftmax along any axis | | `NormKernels` | layer normalisation and RMS normalisation over the last axis, optional scale and bias | | `IntegerKernels` | int32 and int64 Add, Sub, Mul, Div, Mod, Fmod, Min, Max, Neg, Abs; Equal, Less, Greater, LessOrEqual, GreaterOrEqual over int32, int64 and bool; And, Or, Xor, Not on bool; Fill (`ConstantOfShape`) and Range | | `TransposeKernels` | N-d permutation, NCHW↔NHWC, split and merge of attention heads | | `AttentionKernels` | fused scaled dot-product attention with masks, key lengths, causal rule and grouped heads; scalar, cooperative-matrix and Apple simdgroup variants (see [attention.md](https://recursiveai.org/onegpu/docs/attention/)) | | `ConvKernels` | Conv and ConvTranspose (groups, dilation, padding, depthwise, fused bias, residual and activation; float32 and half), see [conv.md](https://recursiveai.org/onegpu/docs/conv/) | | `PoolKernels` | MaxPool, AveragePool (ceil mode, count_include_pad), GlobalAveragePool, GlobalMaxPool | | `ResizeKernels` | Resize, nearest and linear, five coordinate transformations | | `GatherKernels` | ONNX Gather (any axis, negative indices, int32 or int64 indices; the embedding lookup), GatherElements, GatherND (with batch dimensions), over 8-, 16-, 32- and 64-bit payloads | | `ScatterKernels` | ScatterND and ScatterElements with reduction none, add, mul, min or max | | `LayoutKernels` | Slice (including negative steps), Expand, Concat (up to four inputs per pass, more by chaining; inputs may live in the output's buffer when their ranges do not overlap the output's, read in place, or through a scratch copy in a browser); Expand may run in place (below) | | `Int8Gemm` | MatMulInteger (uint8/int8 in any mix, scalar or per-row/column zero points, int32 result), the float-scaled form and the requantized form of QLinearMatMul | | `ConvIntegerKernels` | ConvInteger (uint8/int8 input and weights in any mix, scalar x zero point, scalar or per-channel w zero point, groups, dilation, padding, rank 3 and 4, int32 result) | | `DynamicQuantizeKernels` | DynamicQuantizeLinear (float32 to uint8 with scale and zero point) | | `DequantizeKernels` | DequantizeLinear: uint8, int8 or int32 to float32 or Half, per tensor or per axis, optional zero point; QuantizeLinear: float32 or Half to uint8 or int8, per tensor or per axis | | `RecurrentKernels` | ONNX RNN, GRU and LSTM over float32 or Half, forward, reverse and bidirectional (see [recurrent.md](https://recursiveai.org/onegpu/docs/recurrent/)) | Adding an operation to the element-wise family is a case in `Shaders/Functions.hlsl` and an enum member: the operation is a push constant, not a pipeline. ## Integer and bool tensors - **int32** is one 32-bit word per element. **int64** is two words (low, then high) handled with 32-bit arithmetic, so a device needs no 64-bit integer support; every value from `long.MinValue` to `long.MaxValue` is exact, and so is division (a shift-and-subtract loop runs only when an operand does not fit 32 bits). - **bool** is a `GpuBuffer`: one byte per element, 0 or 1, four to a 32-bit word. It is the same layout as a CPU `bool[]`, so a model's bool tensor uploads as it is. Allocate `GpuKernelBuffers.PaddedBoolLength(n)` elements; a buffer that ends inside a word throws `GpuKernelShapeException`. A thread owns whole output words, so a kernel that writes a bool range at an odd offset keeps the neighbouring bytes of the words it shares, and no atomics are needed. Inputs treat any nonzero byte as true; outputs are 0 or 1. - Arithmetic wraps (two's complement). `Div` truncates toward zero, is 0 for a zero divisor, and the minimum value divided by -1 is itself. `Mod` takes the divisor's sign, `Fmod` the dividend's; both give 0 for a divisor of 0 or -1. These are onemodel's CPU semantics, which the tests check bit for bit, as they do every operation of the family: integer, bool and cast results have no tolerance. - `Range` takes its start and delta as host values (a graph's `Range` inputs are known when a model's shapes are planned) and its length from `IntegerKernels.RangeLength`, which follows the ONNX formula in exact integer arithmetic. The float range is `start + i * delta` as two separately rounded float32 operations. - `Cast` to bool is "not equal to zero" (NaN counts, negative zero does not); from bool it gives 0 or 1. - One variant serves `Fast` and `Deterministic`: integer arithmetic has no order to fix and no atomics are used. - An int64 `Where` moves two words per element and so needs one more dimension than the other types: its shapes may have one fewer dimension after merging. ## Binding caches A family builds a bind group the first time it sees a set of buffers and reuses it on every later call with the same set, so a warm call allocates nothing. Every family constructor takes `maxCachedBindGroups` (default `GpuKernelBindings.DefaultMaxCachedBindGroups`, 1024). When a family has seen more sets than that it drops all of them and starts again, so the bound should exceed the number of distinct buffer sets one run records (a smaller bound is correct, only slower: a dropped group is destroyed after the command buffers that refer to it are submitted, in the browser too): a 72-weight transformer touches 72 sets in `Gemm` per pass, and with a bound of 64 every pass rebuilt them all (72, then 144 bind groups over two passes) where a bound of 72 or more builds 72 once. `ReleaseBindings()` on every family (and `Gemm`) drops the cached groups immediately; call it after disposing buffers that will not be used again, or after replacing a model's buffers, so stale entries do not keep them reachable. The pipelines stay, and the next call with a set of buffers builds its group again. `BindGroupsCreated` on every family (read-only, never reset) counts the groups built so far, so a consumer can record a run twice and check that the second run adds none; for `RecurrentKernels` it includes the input projection's `Gemm`. ## In place Slice and Concat refuse an output that is an input buffer. `Expand` accepts one buffer as both input and output (the way a memory plan places two tensors in one arena) when the element ranges are disjoint: the input's `offset .. offset + count` and the output's `offset .. offset + total` must not overlap, and for `Half` neither may share a 32-bit word with the other. Overlapping ranges throw `GpuKernelShapeException` (operation `Expand`, argument `output`) before anything is recorded. On Vulkan and Metal the same buffer is bound twice in the pass (read-only and read-write); the ranges being disjoint, no thread reads what another writes. A browser refuses one buffer that is both in one dispatch, so there the call expands into a scratch buffer the family keeps (grown on demand) and copies it to the output: two passes instead of one. ## Half precision Half tensors are stored as pairs in 32-bit words and converted in the shader with integer arithmetic, so they need no 16-bit storage feature and run on every backend, including a browser without `shader-f16`. Arithmetic is in single precision and the result is rounded to nearest even once; the rounding is written out in the shader because the drivers' own float-to-half conversion does not promise a rounding mode. - A buffer the kernels touch must end on a word boundary: allocate `GpuKernelBuffers.PaddedHalfLength(n)` halves. Otherwise `GpuKernelShapeException`. - An output may overlap an input only with the same offset and extent. - Softmax and normalisation write whole rows as words, so half tensors need the softmax axis last, an even row length and even element offsets. - Index arithmetic avoids integer division where it can (same-shape and bias-style broadcasts, a single run of reduced dimensions): on Apple GPUs a chain of divisions per element made binary add and reductions two to three times slower, so those cases are compiled as separate variants. - Reduction outputs are written with two integer atomic AND/OR operations on disjoint bits, so neighbouring outputs of different workgroups never clash. There are no floating-point atomics anywhere. ## Tolerances The error measure is max|gpu − reference| / max|reference| over an output, with the reference computed on the CPU in double precision (inputs first rounded to half for half rows). The limits below are what the tests assert; the measured values are far inside them (see the parity table the tests print). | Operation | f32 | f16 (f32 accumulation) | |---|---|---| | Add, Sub, Mul, Min, Max, Where, compare, Cast, Transpose, every `IntegerKernels` operation, Fill, Range (float: bit-exact against `start + i * delta`), Identity, Neg, Abs, Relu, ArgMax/ArgMin | bit-exact (Cast and Transpose also for int and long) | bit-exact vs. `(Half)((float)a op (float)b)` | | Div | 1e-5 (measured 3e-8) | 2e-3 | | Pow, Exp, Log, Tanh, Erf, Sigmoid, Gelu, Silu, Hard*, LeakyRelu, Clip, Sqrt, Reciprocal | 1e-5 | 2e-3 | | Sum, Mean, L1, L2, SumSquare, LogSum | 1e-5 | 2e-3 | | Max, Min | exact | exact | | Prod | max(1e-5, 1e-6·√n) for n factors: every multiplication rounds | 2e-3 | | Softmax, LogSoftmax | 1e-5 | 2e-3 | | Layer and RMS normalisation | 1e-5 | 2e-3 | | RNN, GRU, LSTM (Y, Y_h, Y_c) | 1e-5 | 2e-3 | | QuantizeLinear | bit-exact | bit-exact (Half input widened exactly) | The tanh form of GELU computes its exponential with its own range reduction and polynomial (relative error below 1e-7) rather than the driver's `exp`, whose accuracy is implementation-defined. Erf is the Abramowitz–Stegun 7.1.26 approximation (absolute error below 1.5e-7); the Gelu erf form uses its complementary function so the negative tail keeps its relative accuracy. A float sum of n values has an error of about ε·√n·(Σ|x|/|Σx|) in the worst conditioning: a sum that cancels to almost nothing cannot be accurate in relative terms in any float implementation. The tests use mostly positive data for sums. ## Float32 accumulation order Float32 products and convolutions accumulate k in blocks with block sums added in increasing order (see `gemm.md`), so their error against a double-precision sum grows far more slowly than one chain over k would. Measured scaled error at K = 3072: 1.5e-7 to 3.5e-7 on every device (a single chain gave 1.8e-6 to 4.5e-6). The tests that check it are `AccumulationAccuracyTests`. ## Fast and deterministic The reductions, softmax and normalisation have both modes. - **Fast**: each output (a reduction element, a softmax row, a norm row) is computed by one 256-thread workgroup. Every thread folds a strided slice in index order (reductions add short runs of 32 first, which keeps the rounding error near that of a pairwise sum), and the 256 partial results are combined with the device's subgroup arithmetic when `GpuCapabilities.Compute` reports basic, vote and arithmetic subgroup operations and no subgroup is smaller than 4 lanes; otherwise with the shared-memory tree. `Variant` on each family says which. Constructing a family with `allowSubgroups: false` forces the shared-memory path. - **Deterministic**: the same slices, then a fixed halving tree in shared memory (thread *i* combines with *i + s* for *s* = 128, 64, … 1). No subgroup operation, no floating-point atomic, no dependence on occupancy, scheduling or subgroup size. Two runs on the same device and driver give identical bits, and because nothing depends on the subgroup size the result is also the same for the fast variant without subgroups. - A reduction with fewer than 32 elements per output, or between 32 and 255 along a non-contiguous axis with thousands of outputs, and a softmax row shorter than 16, is done by one thread per output in index order; that path is deterministic by construction and used by both modes. - A softmax row of 16 to 1024 elements gets a power-of-two group of threads (about four elements each), so one workgroup handles 256 / group rows. The statistics are merged by the fixed halving tree in both modes. - A reduction with fewer than 128 outputs and at least 16384 reduced elements (a sum over a whole tensor) runs in two passes: the reduced elements are cut into at most 256 contiguous chunks of at least 4096, a workgroup reduces each chunk to a partial result in a scratch buffer of 128 x 256 floats (`ReduceKernels.ScratchLength`; the family keeps its own unless the caller passes one), and one workgroup per output combines the partials in index order. The cut is a function of the shape, so the deterministic mode stays bit-identical run to run. - Otherwise a workgroup takes each output. The choice depends on the shape only, never on the device. - A workgroup beyond the rounded-up grid (the grid is two-dimensional, so it can hold more workgroups than there are outputs) returns before its first barrier. The workgroup id is uniform, so WGSL's uniformity analysis accepts the barriers that follow. The multi-row softmax keeps a per-thread bounds flag instead, because its rows differ between the threads of one workgroup. ## Gather, scatter, layout and int8 - **8-bit payloads.** `GatherKernels` also copies `byte` (uint8) and `sbyte` (int8) elements, for int8 embedding tables. They are packed four to a 32-bit word like bool tensors (allocate `GpuKernelBuffers.PaddedBoolLength`), one thread owns an output word, and the result is bit-exact. Scatter, concat and the other layout kernels do not take them. - **DequantizeLinear** (`DequantizeKernels`) computes `(float)(x - zeroPoint) * scale` as onemodel's CPU kernel does: the subtraction wraps in 32 bits, the conversion rounds to nearest even, one float32 multiply follows. float32 output is bit-exact with the CPU on every device tested (including int32 inputs near the extremes); Half output rounds that value to nearest even. Per-axis scales and zero points follow the ONNX `axis` attribute; the scale is float32; a zero point has the element type of the input. Blocked quantization and 16-bit inputs are not covered. - **QuantizeLinear** (`DequantizeKernels.EncodeQuantize`) computes `saturate(round_half_even(x / scale) + zeroPoint)` as onemodel's CPU kernel does, bit for bit: the division is the correctly rounded float32 division (done in integer arithmetic, as `DynamicQuantizeKernels` does, because a GPU's own division is not promised to be correctly rounded), the rounding is to nearest even (ONNX operator definition), the zero point is added in float32, the result is clamped to the output type's range, and NaN becomes the lowest value. The output is uint8 (`byte`) or int8 (`sbyte`) and so is the zero point; Half input is widened exactly first. Both modes give the same bits. Tested against the CPU formula on ties, saturation, NaN, infinities, subnormals and five scale sets: 0 mismatching bytes on every device. - **Payloads** are 32-bit (float, int, uint for gather) or 64-bit (long) words, and Half (as 16-bit words). Results are **bit-exact** against the CPU for every payload, since these kernels only move values (scatter reductions on floats do one operation per update, in a fixed order). - **Invalid indices.** A host-side `Validate` helper checks indices before encoding. On the device, an out-of-range index writes zero (gather) or is skipped (scatter), and, when the caller passes a `GpuIndexStatus`, the smallest offending position is recorded; `ThrowIfFailedAsync` raises `GpuKernelIndexException` naming it. Without a status object the zeros are silent. - **Scatter reductions are deterministic in both modes.** Duplicate indices combine strictly in index order: every output element has one owner thread that scans the updates in order (a bucketed scheme, so large scatters still run in parallel). There are no floating-point atomics. `indicesAreDistinct: true` selects a plain parallel path when the caller knows no two updates meet. Ties in min/max keep the existing value, NaN propagates, integer add and mul wrap. Half scatter uses AND/OR on whole words, so it needs the same padding as other Half buffers. uint payloads are not supported for scatter. - **Int8 GEMM** computes `sum (a - zeroA) * (b - zeroB)` in int32 (wrapping), so the result is exact and one variant serves both modes; it matches onemodel's CPU int8 operators bit for bit. The float form is `float(sum + intBias) * scale` followed by a float bias and the int32 form adds its bias to the sum; both are applied as each sum is stored, so a product is one dispatch. The float steps keep the two separate roundings of the unfused chain (the product of the multiply passes through an OR with a bit the host keeps zero, so no target can fuse it with the bias add); the tests compare against the chain on the CPU bit for bit. The requantized form is `clamp(round_half_even(float(sum + bias) * multiplier) + zeroPointY)`; it writes bytes that neighbouring outputs share, so it keeps the sums in a scratch buffer for a second pass. That scratch is the instance's own and grows on demand (a larger product then builds a new bind group), unless the caller places it: `Int8Scratch` (buffer plus element offset) on `EncodeRequantized`, sized with `Int8Gemm.ScratchElements`, or `ReserveScratch` once at session creation. `Int8Gemm.Variant` reports the variant chosen for the device, and `Int8GemmOptions.AllowPackedDot = false` forces the portable one; both give the same bits. A product with more tiles along one axis than a dispatch grid holds (a very wide or tall one: more than 65535 tiles, that is 4.19 million columns at 64-wide tiles) numbers its tiles row-major over x and y; the tile a workgroup owns still depends only on its id, so the bits do not change. - **Packed dot** (`packed-dot-tile64x64x32`): chosen when the device reports an 8-bit dot product (`GpuComputeCapabilities.Int8DotProduct`: Vulkan `shaderIntegerDotProduct`, the browser's `packed_4x8_integer_dot_product`). Four byte pairs are multiplied and summed per instruction (`dot4add_i8packed`). That instruction takes signed bytes only, so unsigned operands are staged as `u - 128` and the zero points move by the same amount; the zero points are then applied after the loop from byte sums collected while staging: `sum (s - z')(t - w') = sum s*t - w' * sum s - z' * sum t + K * z' * w'`, exact modulo 2^32 like the sums themselves (Jacob et al., 2018, section 2.3 rearranges the same product). A byte sum is only collected when the other operand's shifted zero point can be non-zero. On a discrete GPU, products with fewer than 32 tiles of 64x64 run on 32x32 tiles so that the 128-row products of small models still fill the device (an integrated GPU took less time on the big tile everywhere). - **Portable** (`portable-tile64x64x16`): runs everywhere. Rows are staged a word at a time when the operand's offset, stride and leading dimension are multiples of four, and the products are formed in float, which is exact here (two values in -255..255 multiply to under 65025, so 256 of them sum to under 2^24); the running sums fold into int32 every 128 values of k. Metal has no packed dot product, so Apple GPUs run this one. - **DynamicQuantizeLinear** runs in two passes (partial minimum and maximum per workgroup; then every workgroup combines the partials, the first stores scale and zero point, each quantizes its slice) and divides in integer arithmetic (a long division of the mantissas with round-to-nearest-even), because the targets' float division is not correctly rounded (Vulkan allows a few ulp, Metal compiles with fast math) and the scale and zero point must equal the CPU's bit for bit. An int8 linear layer is thereby three dispatches (two to quantize, one product) instead of six. - **Long submissions on Metal.** The system aborts a command buffer that runs too long (`kIOGPUCommandBufferCallbackErrorImpactingInteractivity`). Encoders accept `ReportWork(operations)` and `SplitSubmission()`; Metal splits the command buffer at pass boundaries once a budget (`GpuComputeCapabilities.SubmissionWorkBudget`, 2^33 operations) is reached. `Gemm` cuts a product larger than the budget into row or batch pieces with the same accumulation order, so deterministic results are unchanged (`GemmOptions.AllowSplitting`). A command buffer that still fails surfaces as `GpuCommandFailedException` (`Failure` says watchdog, out of memory or invalid resource) from the next submit, wait or readback; never as a silent wrong result. The browser backend has no budget. ## Special values NaN and infinity follow IEEE 754 on Vulkan and Metal: Min, Max and the reductions Max/Min propagate NaN; ArgMax/ArgMin treat NaN as the extreme value (first NaN wins); softmax of a row with a NaN, a +infinity, or only −infinity gives NaN for the row, while −infinity beside finite values gives 0 (so additive masks work). The kernels test NaN with bit patterns, because Metal compiles with fast-math. **In a browser, WGSL does not require NaN or infinity to behave** (implementations may assume they never occur), so those cases are not tested on WebGPU and masks should use a large negative finite value there. ## Limits Element counts are below 2³¹ and every index is 32-bit. Broadcasting merges neighbouring dimensions that every operand walks contiguously, then needs at most 6; a reduction needs at most 4 runs of kept and 4 runs of reduced dimensions after merging; a transpose at most 6 unmergeable dimensions. Workgroups are 256 threads, which every Vulkan, Metal and WebGPU device offers; dispatches are two-dimensional so more than 65535 workgroups work everywhere. # Matrix multiplication (`Gemm`) `RecursiveAI.OneGpu.Kernels.Gemm` computes `C = alpha * op(A) * op(B) + beta * C` for `Batch` independent products. `op(X)` is `X` or its transpose. A and B are float32, float16 or bfloat16; products and sums are float32 and C is float32. Matrices are dense row-major with a leading dimension; sizes, strides and offsets count elements. The class is created once per element type and mode, picks its variant from the device's capabilities at creation, and records into a caller's `GpuCommandEncoder`. ## Variants | Variant (`Gemm.Variant`) | Used when | Notes | |---|---|---| | `tiled-64x64x16` | everything else, and always in `Deterministic` mode | HLSL through the shader compiler: SPIR-V, MSL and WGSL. | | `cooperative-matrix-16x16x16-tile64x64x16-8sg` (`-4sg`) | Vulkan, `Fast`, f16 or bf16 inputs; subgroups of 32 (fixed, or selectable and pinned to 32); pipelines with full subgroups; the device lists a 16x16x16 combination of the input type with float32 accumulation; the kernel's shader features (16-bit storage, `VK_KHR_cooperative_matrix`, the Vulkan memory model, and for bf16 `VK_KHR_shader_bfloat16`) | HLSL `CooperativeMatrix` through the shader compiler. Eight subgroups per workgroup; four if the device allows fewer than 256 invocations or too little workgroup memory. | | `simdgroup-matrix-16x16x16-tile64x64x16-8sg` (`-4sg`) | Metal, `Fast`, the device lists an 8x8x8 combination of the input type with float32 accumulation, 32-wide simdgroups | The same HLSL; the shader compiler lowers a 16x16 matrix to 2x2 `simdgroup_matrix` 8x8 tiles. Float32 inputs use it too. The eight-subgroup float32 shape declares 18944 bytes of threadgroup memory; the constructor takes the shape from the length the compiled pipeline reports (`GpuComputePipeline.ReportedWorkgroupMemoryBytes`), so with Metal shader validation on, which doubles the reported figure to 37888 against the 32768-byte limit, it falls back to four simdgroups, and without validation it keeps eight. A three-component vector in workgroup memory takes 16 bytes in MSL, which a size formula would also miss. | On Metal the float32 matrix kernel reads its fragments straight from the operands for every tile that lies inside C and every whole 16 steps of k, and stages only edge tiles and the ragged k tail (`GEMM_DIRECT`), with four subgroups: 1024 x 3072 x 768 ran at 1488 GFLOP/s against 1065 for the staged eight-subgroup kernel, and 1024 x 768 x 3072 at 1470 against 869 (Apple M2, host clock, indicative). `GemmOptions.AllowMatrixHardware = false` forces the portable variant. WebGPU has no matrix hardware in the browser today and always runs the portable variant, whose half-precision inputs are read as pairs inside 32-bit words, so it needs no 16-bit storage feature and runs in Chrome without `shader-f16`. **Which devices take the matrix kernel.** The choice is made once, in the constructor, from `GpuCapabilities.Compute` and the kernel's own reflection (`EntryLayout.RequiredFeatures` against `ShaderFeatures`); nothing is guessed from the vendor. The kernel is written for 32-wide subgroups: its pipeline is created with full subgroups and, where the device can run several widths, with the subgroup size pinned to 32 (`ComputePipelineDesc.RequireFullSubgroups`, `RequiredSubgroupSize`). A device that only has 64-wide subgroups, such as an AMD GPU in wave64 mode with no wave32 option, does not qualify and runs the portable kernel; a wave64 variant has not been written because no device that offers one was available to test it on. The AMD GPU in the RADV test machine (Raphael iGPU, subgroups 32 to 64, size selectable) reports no cooperative-matrix shapes at all, so it takes the portable kernel as well, as does lavapipe. ### Portable kernel One 64 x 64 tile of C per workgroup of 256 invocations; the k loop runs in steps of 16 through two workgroup-memory tiles (Goto and van de Geijn blocking); every invocation owns a 4 x 4 register tile with interleaved rows and columns. Tiles are staged with bounds checks, so any M, N, K, leading dimension and offset works. A grid of 64 x 64 tiles has `ceil(M/64) * ceil(N/64) * Batch` workgroups. ### Matrix-hardware kernels The same 64 x 64 output tile and 16-wide k step as the portable kernel. Eight subgroups (2 x 4) each own 32 x 16 of C as two 16 x 16 float32 accumulators; the four-subgroup shape (2 x 2) owns 32 x 32 as four. Each k step copies the A and B tiles from the operand buffers, in their own element type, into padded workgroup memory (this is where the bounds, the transposes and odd leading dimensions are handled), then every subgroup loads its fragments and multiplies them with cooperative matrices. The accumulators leave through workgroup memory so alpha, beta and the edges are plain code. The copy involves no conversion: reading bfloat16 through float32 first cost a third of the throughput. The 16-bit elements are read from storage buffers as 16-bit values, which needs 16-bit storage access; every device with cooperative matrices has it. ## Deterministic mode `GpuKernelMode.Deterministic` always uses the portable kernel. Each element of C is summed by one invocation, in increasing k, in float32 with fused multiply-add, or, where the grid is small (see "Small grids and few rows"), as range sums added in range order, in both modes; there is no atomic, no subgroup arithmetic and no dependence on how workgroups are scheduled, so two runs on one device and driver give the same bits. Different devices and drivers may differ in the last bits (fused multiply-add contraction and rounding of the compiler's output are theirs). `Fast` may use another summation order and the matrix-hardware variants, so `Fast` and `Deterministic` results differ by rounding. ## Small grids and few rows (portable kernel, Vulkan and WebGPU) Two cases leave a large GPU idle with the 64 x 64 tile: a small grid over a long k, and a product with only a few rows. Both are chosen from the shape alone and serve `Fast` and `Deterministic` alike, because both fix the order in which an element's terms are added. - **Split k.** With fewer than 384 workgroups of 64 x 64 (a 128 x 768 product has 24) and at least 128 terms of k, k is cut into up to 16 ranges of a multiple of 32 terms (at least 64 terms each), aiming at about 384 workgroups. Each range writes its raw sums to its own plane of an instance-owned 16 MiB float32 scratch buffer (allocated on first use; fewer ranges when the planes would not fit), and one more pass adds the planes in range order, then applies alpha, beta, the bias and the activation. A product cut into row or batch pieces for the submission budget uses the ranges of the whole product, so its bits do not change. Not used on matrix hardware, whose kernels are untouched. - **Few rows.** A float32 product with at most four rows of A (decoding one to four tokens), A not transposed, runs a kernel that streams B once: for B stored K x N a workgroup owns 16 columns (32 from N = 8192, so a warp reads 128 bytes of a row) and its 256 threads split k into 16 (8) interleaved slices; for B stored N x K a workgroup owns 8 columns and 32 threads per column read consecutive k. Every thread chains its terms in increasing k, and one thread adds the slice (lane) sums in order through workgroup memory. Not on Metal. Float16 and bfloat16 inputs keep the tile. Measured on the RTX 4070 (GPU time from timestamps, float32, indicative; before is the 64 x 64 tile alone): | Product (M x N x K) | Before | After | | |---|---|---|---| | 128 x 768 x 768 (BERT-base projection, batch 1) | 58 us | 29 us | 2.0x | | 128 x 2304 x 768 | 94 us | 76 us | 1.2x | | 128 x 3072 x 768 (feed-forward up) | 135 us | 100 us | 1.35x | | 128 x 768 x 3072 (feed-forward down) | 222 us | 100 us | 2.2x | | 1024 x 768 x 3072 (BERT batch 8, feed-forward down) | 925 us | 758 us | 1.2x | | 1 x 768 x 768 | 57 us | 5.9 us | 9.7x | | 1 x 3072 x 768 | 93 us | 11 us | 8.4x | | 1 x 768 x 3072 | 221 us | 15 us | 15x | | 1 x 51865 x 768, B stored N x K (decoder logits) | 771 us | 330 us | 2.3x (159 MB in 330 us: 480 GB/s) | | 1 x 51865 x 768, B stored K x N | 768 us | 337 us | 2.3x | | 4 x 3072 x 768 | 93 us | 11 us | 8.6x | After a split the 128-row products run at 6.0 TFLOP/s, the tile's rate on large shapes (5.1-6.9), so the next step for them is a float32 tile with a higher rate, not more splitting. ## Accumulation order Rounding in one chain over k grows roughly in proportion to k, because every addition rounds a partial sum that keeps growing (Higham, *Accuracy and Stability of Numerical Algorithms*, chapter 4). The float32 kernels sum in blocks and add the block sums, so each rounding happens on a smaller partial sum. The order depends only on the shape and the variant: no atomics, no subgroup operations, nothing that varies with scheduling, so two runs on one device and driver are bit-identical in both modes. - Portable kernel (all element types, and the convolution kernel that shares its tile): the 16 products of one staged step chain in a register; two staged steps (32 terms) form a block sum; four block sums (128 terms) form a group sum; group sums join the total in increasing k. Float16 and bfloat16 inputs read through the same kernel use the same order. - Float32 matrix hardware (Apple, and the implicit-GEMM convolution on it): each subgroup keeps a second accumulator that takes eight staged steps (128 terms) and then joins the total, in increasing k. - Float16 and bfloat16 on matrix hardware keep a single chain: the hardware sums each 16-term step before it reaches the accumulator, the inputs' own rounding dominates the error at these types' tolerances, and the second accumulator cost about half the throughput on NVIDIA. - Convolution split along k (Apple small grids) adds its partial results in range order, as before. ## Bias and activation epilogue `Encode` takes an optional `GemmEpilogue`: a float32 bias vector (element `j` is added to column `j` of every row of every product in the batch, from an element offset) and an activation (`Relu`, `GeluErf`, `GeluTanh`), applied to each result element after `Alpha * product + Beta * C`, in the same dispatch and by every variant (portable, cooperative-matrix, simdgroup-matrix). The activations are the functions `ElementwiseKernels` runs (`GpuUnaryOp.Relu`, `GeluErf`, `GeluTanh`), and the bias is one float32 add, so the fused result is the result of the product followed by those kernels: the tests compare the two bit for bit (every dtype, variant, mode, activation, with and without bias) and the comparison is exact on every device tested. Where a compiler fuses `Alpha * product` with the bias add into one fused multiply-add, the last bit can differ for an `Alpha` that is neither 1 nor a power of two; none of the devices tested did. Bias and activation are push constants plus one more binding, not pipeline variants, so the deterministic variant stays deterministic (the epilogue is per element). The bias is float32 whatever the input type, because C is float32. ## Accuracy Metric: `max |gpu - ref| / max |ref|` over the output, with the reference computed in double from the inputs after rounding to the input type. The tests report it together with the largest element-wise relative error over elements whose reference is above 1e-3 of the largest (smaller ones are cancellation noise). | Inputs | Bound | |---|---| | float32 | 1e-5 | | float16 with float32 accumulation | 2e-3 | | bfloat16 with float32 accumulation | 8e-3 | All variants measured far below these, because the inputs are exactly the rounded values and only the float32 accumulation order differs. The float32 accumulation order is described under "Accumulation order" below; with it the scaled error against a double-precision sum is about 1.5e-7 to 4e-7 at K = 768 to 4096 (a single chain over k gave 1e-6 to 4.5e-6 there). Elements outside the written region of C are checked to be untouched. With `Beta = 0` C is never read. ## Limits and errors - Every shape, stride, offset and buffer range is checked in 64-bit arithmetic before anything is recorded; a bad one throws `GpuKernelShapeException` naming the operation and the dimension. The grid must fit `MaxWorkgroupCount` per axis. - Each operand must fit one storage-buffer binding (`MaxStorageBufferBindingBytes`); a larger one throws `GpuUnsupportedException`. Software Vulkan drivers report 128 MB, which a 51865 x 768 float32 matrix exceeds. - Buffers made by `CreateBuffer` are padded to a whole 32-bit word, so a 16-bit element at the end of an odd-length buffer can be read as part of a word. - Recording a call allocates nothing once the bind group for its three buffers is cached (64 triples; `ReleaseBindings` drops them). - A batch with `StrideB = 0` shares one B (a weight matrix) across the batch; `StrideC` must keep the outputs apart. ## Throughput (indicative) GPU time from timestamp queries (median of 30 dispatches) on the RTX 4070, driver 595; host time between submit and idle on the Apple M2 (Metal here reports no timestamp queries). Peaks used: RTX 4070 float32 29.2 TFLOP/s (5888 lanes x 2 x 2.475 GHz); float16 tensor with float32 accumulation about 58 TFLOP/s dense (derived from the Ada figures, so approximate); M2 float32 3.6 TFLOP/s (Apple's figure for the 10-core GPU). These are indicative numbers, measured while other jobs shared the machines. | Device, variant, inputs | Large shapes (GFLOP/s) | Fraction of peak | |---|---|---| | RTX 4070, tiled, float32 | 5100 to 6900 (square 4096: 5500) | 18 to 24 % | | RTX 4070, cooperative matrix, float16, 8 subgroups | 25500 to 34200 (square 4096: 36500) | 44 to 59 % (square: 63 %) | | RTX 4070, cooperative matrix, bfloat16, 8 subgroups | 25100 to 33900 (square 4096: 36300) | 43 to 58 % (square: 63 %) | | RTX 4070, cooperative matrix, float16 or bfloat16, 4 subgroups | 19800 to 29700 (square 4096: 30900) | 34 to 51 % | | Apple M2, simdgroup matrix, float16, 8 simdgroups | 436 to 456 (square 4096: 402 to 407) | 12 % | | Apple M2, simdgroup matrix, bfloat16, 8 simdgroups | 419 to 440 (square 4096: 406 to 435) | 12 % | | Apple M2, simdgroup matrix, float32, 4 simdgroups | 308 to 358 (square 4096: 328) | 9 to 10 % | | Apple M2, tiled | 100 to 139 | 3 to 4 % | The Apple rows were measured on a machine that other jobs were using, so treat them as lower bounds. The float32 matrix kernel on the M2 is about 20 % slower than the 8x8 kernel written by hand that it replaced (308 to 358 against 398 to 452), because the Metal compiler counts the 8-simdgroup float32 shape at twice its declared threadgroup memory and so it keeps four simdgroups; float16 is 2 to 9 % below the old kernel and bfloat16 runs at three to five times the speed (the old kernel converted each element). On the RTX 4070 the float16 kernel is 1.25 to 2.0 times the hand-written SPIR-V it replaced and bfloat16 is 1.3 to 2.2 times (the smallest gain is the batched 128 x 128 scores product, the largest the square 4096 product), because the operands are now copied into workgroup memory without a detour through float32 and eight subgroups hide more latency than four. Single-row products (decode, M = 1) are memory bound. On Vulkan and WebGPU, float32 products of at most four rows have their own kernel (see "Small grids and few rows"): 242 GFLOP/s for the vocabulary projection on the RTX 4070 at 480 GB/s of weights, against 103 with the tile. On the M2 they still use one row of a 64-row tile (2 to 9 GFLOP/s); a few-rows kernel and larger tiles for the Apple GPU are the next tuning steps. ## 8-bit products on Apple matrix hardware `Int8Gemm` on Metal stages each tile as float16 after subtracting the zero points (a value in -255..255 is exact in float16), multiplies on the simdgroup matrices (products exact in float32) and adds float32 sums of at most 2^24 in magnitude. A chunk of 16 steps of 16 (256 terms) is exact for two operands up to 255 in magnitude; 32 or 64 steps when an operand is a plain int8. After each chunk the sums become int32 and add into int32 registers, which wrap like the definition. The result is the same integer for every summation order, so one kernel serves both modes. Tests include K = 4096 at the extremes. `Int8GemmSkinny` takes M <= 16 with a non-transposed B (decoder steps): 32 columns x 4 rows per workgroup, k split 16 ways, integer sums. Indicative on an M2 (GPU timestamps): the matrix tile reaches 0.67-0.77 TOPS at BERT, MiniLM and Whisper-encoder shapes against 0.52-0.58 for the portable tile (1.3x, staging bound); decoder-step products with M = 1 take 28 us (projection), 80 us (ffn up), 97 us (ffn down) and 2.97 ms (vocabulary) against 271, 619, 1099 and 10020 us. A transposed B with M = 1 still takes the portable path. # Attention `AttentionKernels` computes `O = softmax(Scale · Q Kᵀ + Mask) V` for every batch entry and head in one kernel. The score matrix is never written to memory: a workgroup keeps a block of 64 query rows and walks over the keys one tile at a time, keeping for each row a running maximum, a running sum of `exp(s − max)` and the accumulated `exp(s − max) · V`, and rescales them whenever a tile raises the maximum (Dao et al., *FlashAttention*, 2022, for the tiling; Milakov and Gimelshein, *Online normalizer calculation for softmax*, 2018, for the normaliser). ```csharp using var attention = new AttentionKernels(device); var p = new AttentionProblem(batch, heads, qLen, kLen, headDim) { Q = GpuAttentionLayout.PositionsFirst(heads, qLen, headDim), // [batch, qLen, heads, headDim], as a projection leaves it K = GpuAttentionLayout.PositionsFirst(heads, kLen, headDim), V = GpuAttentionLayout.PositionsFirst(heads, kLen, headDim), O = GpuAttentionLayout.PositionsFirst(heads, qLen, headDim), Mask = GpuAttentionLayout.ForMask([batch, 1, 1, kLen], batch, heads, qLen, kLen), // optional additive mask }; var encoder = device.BeginCommands(); attention.Encode(encoder, p, q, k, v, o, mask: maskBuffer, mode: GpuKernelMode.Deterministic); device.Submit(encoder); ``` ## What it covers | | | |---|---| | Element types | float32; `Half` Q, K, V and O with float32 accumulation; bfloat16 (`ushort` bit patterns, `EncodeBFloat16`) on matrix hardware only, head dimension up to 128 (the mask is always float32) | | Layouts | any `GpuAttentionLayout` (offset, batch, head and row strides; elements of a row contiguous): head-first, position-first, slices of a packed QKV projection | | Lengths | query and key/value lengths differ (cross-attention, decoder steps with a key/value cache) | | Heads | equal query and key/value heads, or fewer key/value heads (grouped and multi-query attention) | | Head dimension | 1 to 256; 32 and 64 are tuned, others are padded up to 32, 64, 128 or 256 (above 64 the registers spill: correct, slow). Half needs it even | | Masks | additive float mask broadcast over batch, heads and rows (`-∞` hides a key); per-batch key lengths (`GpuBuffer`); causal with an offset (default `KLen − QLen`, the cached-decoder case) | ## Special values - A key hidden by `-∞` in the mask, by its batch entry's key length or by the causal rule has weight 0. - A query row with no visible key gives **NaN** in every output (what `0 / 0` gives, and what a plain softmax over an all-`-∞` row gives), or zeros with `MaskedRows = GpuAttentionMaskedRows.Zero` (a row that attends to nothing contributes nothing). `KLen = 0` is such a row. - A row whose score is NaN or `+∞` gives NaN. V is not inspected. - In a browser, WGSL does not promise infinities and NaN behave, so use a large finite negative mask value there (`-1e30`); the tests do. ## Variants and ordering Keys are folded into each query row one at a time, in increasing key order, with the online normaliser, from key and value tiles held in workgroup memory. Two kernels share that structure (`AttentionKernels.Variant`, chosen once in the constructor): - **`rows`**: one thread owns a query row (64 rows per workgroup). Needs nothing from the device. - **`pairs`**: two neighbouring threads own a row, each the half of the head dimension, and add their halves of a dot product after one subgroup shuffle (32 rows per workgroup). This halves the registers per thread. Chosen automatically on Apple GPUs, which spill past about 128 registers per thread (a 64-wide head needs 128 for the query and the accumulator); on the RTX 4070 it took longer for large problems and less time for small ones, so `Auto` picks `rows` there. `GpuAttentionVariant` forces either; `allowSubgroups: false` rules out `pairs`. The compiler keeps loops as loops and a thread-local array indexed by a loop variable can live in slow memory (Apple's compiler does that), so the SPIR-V and Metal builds write every access to the query and accumulator registers as a separate statement with a literal index (macros in `Attention.hlsl`). The WebGPU build for the software adapter (`AttentionLoops.hlsl`) keeps plain loops instead: its compiler does not finish compiling the unrolled form. A browser on a real GPU gets the unrolled form like the other targets (see *In a browser* below). ### In a browser Chrome's WGSL compiler did not keep a loop-indexed private array in registers: with plain loops the fused kernel was 3 to 18 times slower than GEMM + softmax + GEMM on an RTX 4070 (Chrome 155, wall time per call from a batch of calls closed by a one-element readback; indicative). Two changes closed the gap, and neither moves a bit on the other targets: - A WebGPU device on a hardware adapter uses the unrolled statements (WGSL is generated from `Attention.hlsl`); a software adapter (SwiftShader) keeps `AttentionLoops.hlsl`. - Key and value tiles are read as `float4` (`VEC4` build). WGSL checks every dynamic index of a workgroup array, so reading four floats per access also removes three of four checks, and the unrolled form's index arithmetic is the same either way. The sums keep the element-wise order (each group of four adds its products in increasing element order), so the results equal the scalar form's bit for bit. Apple GPUs ran the `float4` form 17 to 25% slower (rows variant, M2) and keep the scalar one. | Shape (12 heads) | fused, plain loops | fused, unrolled + `float4` | GEMM + softmax + GEMM | |---|---|---|---| | MiniLM B1 S128 (head 32) | 0.47 ms | 0.14 | 0.16 | | MiniLM B8 S128 | 1.12 | 0.27 | 0.30 | | BERT B1 S128 | 1.15 | 0.14 | 0.15 | | BERT B1 S512 | 5.14 | 0.38 | 0.40 | | BERT B8 S128 | 2.49 | 0.27 | 0.28 | | BERT B8 S512 | 25.97 | 1.23 | 1.93 | | Whisper encoder, 1500 keys | 41.20 | 1.93 | 2.30 | | Whisper decoder, 448 cached keys | 1.32 | 0.14 | 0.15 | | Whisper cross-attention, 1 query | 2.31 | 0.14 | 0.31 | The 0.14 ms rows are the floor of one submit and readback in the browser, not kernel time. The same kernel in Vulkan on that card takes 1.09 ms for the Whisper encoder (GPU timestamps), so the browser form is within 1.8x of native (wall time, which includes the readback). `AttentionBrowserTiming` reruns the table in a browser on a hardware adapter. ### Long calls A single submission that keeps an Apple GPU busy for a long time (the one-thread-per-row kernel on a batch of 8 Whisper encoder inputs was the example) is stopped by the system and leaves the output half written. `Encode` therefore cuts a call whose estimated work exceeds the device's `SubmissionWorkBudget` (`submissionWorkBudget` in the constructor overrides it; 0 never cuts) into dispatches of whole workgroups and tells the encoder the work after each (`ReportWork`), so the encoder can end the submission between them. Every workgroup computes its own rows from the same inputs in the same order, so a cut call equals the whole call bit for bit; the tests compare them on every device. The estimate is `4 · batch · heads · QLen · KLen · HeadDim` operations, times 1.5 for the matrix kernel (it multiplies `Q Kᵀ` twice) and times 4 for the scalar kernels (several times slower per operation, measured on the M2). ### Accuracy at long key lengths Adding one product per key into a single running sum lets the rounding error grow with the number of keys: at 4096 keys with moderately spread scores the plain sequential form was several times less accurate than GEMM + softmax + GEMM (whose sums are shorter or tree-shaped). For head dimensions 32 and 64 the scalar kernels therefore sum blockwise: the products of one tile of keys go into a second accumulator that starts at zero, and each finished tile is added to the main accumulator after one rescale by `exp(old max - max)`. Every sum is over one tile (32 keys) or over the tiles, in a fixed order, so determinism is unchanged. The price is `PD` more registers per thread (a batch of 8 BERT-base inputs ran about 30% slower on the RTX 4070), so only workgroups that walk 640 keys or more use it (the cut ranges of a split call are usually shorter), and head dimensions 128 and 256 keep the sequential form (they spill already). Rounding the float32 scores themselves is a floor both paths share: with scores spread over a standard deviation of about 12 or more, both the fused and the unfused path sit near 1e-5 whatever the summation. `Fast` and `Deterministic` run the same kernel: every sum has an order the code fixes, there are no atomics and no subgroup arithmetic, and the one shuffle adds two numbers (which does not depend on operand order). Two runs on one device and driver give identical bits. `rows` and `pairs` add a dot product in a different order, so they may differ from each other, and devices from each other, in the last bits. When a problem has too few query blocks to fill a GPU, the keys are cut into at most 16 contiguous ranges of a multiple of 32 keys (at least 128 keys each, enough ranges for about 256 workgroups); each range stores its partial (maximum, sum, accumulator) in a scratch buffer and a second pass merges them in range order. The cut depends only on the shape. `ScratchLength(problem)` says how many floats are needed (0 without a cut); pass a buffer as `scratch`, or leave it null and the family keeps one of its own and releases it on `Dispose`. ### Matrix hardware (`cooperative-matrix`) For half and bfloat16 tensors with a head dimension up to 128, `Auto` picks a kernel that runs `Q Kᵀ` and `P V` as 16 x 16 x 16 cooperative matrices (`VK_KHR_cooperative_matrix` on Vulkan, simdgroup matrices on Metal), when the device lists such matrices of the element type with a float32 accumulator, runs 32-wide subgroups (pinned where the device can run several widths) with full subgroups, and offers the kernel's shader features. Everything else (float32, head dimensions 129 to 256, devices without matrix hardware such as lavapipe and the RADV iGPU, WebGPU) keeps `rows` or `pairs`; `GpuAttentionVariant.Cooperative` forces the matrix kernel (and throws where it cannot run) and `allowMatrixHardware: false` rules it out. `Variant` names the kernel used for half tensors, `FloatVariant` the one for float32. A workgroup owns 64 query rows as four subgroups of 16 rows and visits the keys in tiles of 32 (16 for 128-wide heads), staged through workgroup memory (at most 12.8 KB, so it also fits a Metal limit that shader validation counts twice). The elements of a cooperative matrix have no specified layout, so an accumulator cannot be scaled per row without a round trip through workgroup memory; the kernel avoids the rescaling of the online normaliser instead. A first pass over the keys computes only the scores and finds the exact maximum of every row; a second pass recomputes them, forms `p = exp(score − max)` and accumulates `P V` in matrix registers. The price is a second `Q Kᵀ`, 1.5x the matrix work. The probabilities are rounded to the element type before the product with `V`, and the normaliser sums the rounded values, so numerator and denominator see the same `p`. Scores and probabilities pass between the matrix units and the lanes through workgroup memory, where lane `l` of a subgroup owns row `l & 15` and every other column of each 16-wide block; the two lanes of a row meet once per pass in a shuffle that maximises or adds two numbers (operand order does not matter). The key-range split and the merge pass are those of the scalar kernels. This kernel serves `Fast` and `Deterministic` alike: every order is fixed by the code (tiles in increasing key order, the hardware's matrix product depends only on its operands), there are no atomics and no subgroup arithmetic, and two runs on one device and driver give identical bits (checked per device by the tests). Devices differ from each other in the last bits, and so do the matrix and scalar kernels. ### Apple simdgroup kernel (`AttentionSimd.hlsl`, Metal only) On Apple GPUs every element type (float32 with head dimension up to 64, half, bfloat16; head dimensions 32, 64 and 128) runs a single-pass kernel on 8 x 8 simdgroup matrices, in place of `pairs`/`rows` (those stay as the fallback and for `GpuAttentionVariant` overrides). `Variant` and `FloatVariant` name the kernel. The scalar kernels were 10-20x slower than GEMM + softmax + GEMM here; this one is at or below the unfused chain at every reference shape (table below). - **Scores transposed.** `Sᵀ = K Qᵀ`, so a row of the softmax is a column of the matrix and each lane owns the same row across blocks; no per-row scaling of an accumulator through memory is needed unless the maximum changes. - **Reference maximum.** Each row keeps a reference `m` raised only when a block's maximum exceeds it by more than 8 (`REF_SLACK`): `exp(score − m)` is then at most `e⁸`, far from overflow, and the accumulators are rescaled by `exp(old − new)` through workgroup memory only on a raise (the decision is made per workgroup from flags, so the barriers stay uniform). The first block sets `m` with no rescale. Whether a rescale happens depends only on the data, never on timing. - **Query-direct variant.** For 16 or more queries the 4 simdgroups of a workgroup (64 rows) reload their Q fragments from the tensor at each step instead of keeping them in registers (the last 10-15 %); the last 16-row block slides back so no row is written twice. Decoder steps and very short queries keep Q in registers, one simdgroup per 16 rows. - K and V fragments load straight from the buffers when the key block is full and the head dimension is a multiple of 16; otherwise through a staged tile. Indicative M2 wall time, ms, validation off, 12 heads of 64 (Whisper encoder) or BERT/MiniLM heads: unfused f32 GEMM + softmax + GEMM against the fused kernel before (`pairs`) and after. | Shape | unfused f32 | fused f32 before | fused f32 after | fused f16 after | |---|---|---|---|---| | Whisper encoder, 1500 keys | 13.7 | 26.7-32 | 10.3 | 8.9 | | BERT-base B8 S512 | 11.5 | 18-19 | 9.3 | 8.1 | | BERT-base B1 S512 | 2.0 | | 1.4 | 1.2 | | MiniLM B8 S128 | 1.03 | | 0.47 | 0.42 | | Whisper decoder, 448 cached keys | 0.24 | | 0.21 | 0.20 | | Whisper cross-attention, 1 query | 0.75 | | 0.30 | 0.17 | Accuracy (scaled, double reference): float32 1.5e-6 on the model shapes and 4.5e-6 on scores that keep rising or are wide (the same as the scalar kernels there), half 5.0e-4, bfloat16 4.4e-3; fused against unfused float32 at most 9.7e-7. Two runs are bit-identical. As with the other kernels, `Fast` and `Deterministic` run the same code. The 8x8 product is a fixed function of its operands and each output tile is summed by one simdgroup in key order, so it counts as a fixed order; the scores and the O rescale use no atomics or subgroup arithmetic. ## Tolerances Scaled error (max |gpu − reference| / max |reference| per output, reference in double on the CPU, inputs rounded to half first for half rows): float32 ≤ 1e-5, half with float32 accumulation ≤ 2e-3, bfloat16 with float32 accumulation ≤ 8e-3 (the probabilities and the output are rounded to bfloat16's eight bits of precision, which costs about 4e-3). The tests assert these on the shapes of the reference models (MiniLM, BERT-base, Whisper-small encoder, decoder self- and cross-attention) and on tile edges, remainders, size 1, non-power-of-two sizes, masks with `-∞`, fully masked rows, key lengths of 0, grouped heads and odd head dimensions; they also compare the fused kernel with GEMM + add + softmax + GEMM. # Convolution, pooling and resize `ConvKernels`, `PoolKernels` and `ResizeKernels` in `RecursiveAI.OneGpu.Kernels` cover ONNX `Conv`, `ConvTranspose`, `MaxPool`, `AveragePool`, `GlobalAveragePool`, `GlobalMaxPool` and `Resize` for the vision models (ResNet-50, MobileNetV3, DBNet, a recurrent-recognizer's convolutional stem). They follow the binding shape in [kernels.md](https://recursiveai.org/onegpu/docs/kernels/): one class per family, created from a `GpuDevice`, operands as `GpuBuffer` plus element offsets, shape errors as `GpuKernelShapeException`. ## Layout and operands - Tensors are **NCHW**, weights **OIHW** (`ConvTranspose`: `[Cin, Cout/groups, KH, KW]`). One-dimensional operators are two-dimensional with height 1. NCHW is what ONNX graphs carry, so a graph needs no layout pass; NHWC would only pay with one. - Element types: `float` and `Half` (half in and out, float32 accumulation). Pooling and resize also move any 32-bit or 16-bit element bit-for-bit where the operation copies. - A problem struct (`ConvProblem`, `ConvTransposeProblem`, `PoolProblem`, `ResizeProblem`) carries the shape, explicit pads, dilation, groups, the fused activation and element offsets. `WithAutoPad` and `GpuWindow.OutputExtent` / `TransposeOutputExtent` turn ONNX `auto_pad`, `output_padding` and `output_shape` into pads and output sizes. Sizes are checked with `checked` arithmetic and the device's limits before anything is recorded. - Fused into the last pass, in this order, each step rounded in float32: sum, bias, residual addition, activation. Activations fused into the matrix kernel: identity, tanh, sigmoid, relu, silu, hard-sigmoid, hard-swish, clip, leaky-relu; any other `GpuUnaryOp` runs as one extra element-wise pass. - Scratch (the column matrix of the unfold variants and, for half, float32 accumulators) is supplied by the caller as a buffer **and an element offset** (`columns`/`columnsOffset`, `accumulators`/`accumulatorsOffset`, `EncodeTranspose` likewise) or kept by the instance up to 64 MiB. `ConvKernels.Scratch(problem, half)` says how much this device needs for a shape (nothing for the variants that gather windows while they multiply), for the whole batch and for one image; a smaller buffer makes the batch run in chunks of whole images, which changes no bit of the result. The scratch may live in the buffer that holds the operands (an arena) as long as its range overlaps no operand's (`GpuKernelShapeException` otherwise). A browser cannot read and write one buffer in a dispatch, so there the scratch needs a buffer of its own, and so do the output and the inputs. ## Variants `ConvKernels.DescribeVariant(problem, half)` names what a device runs. | Layer class | Variant | |---|---| | depthwise (one input channel per group, any channel multiplier) | direct kernel, one thread per output element, window order fixed | | every other convolution on the portable matrix kernel (float32 everywhere; half on devices without matrix hardware) | **implicit GEMM**: a 64x64x16 tile kernel that gathers the input window while it loads the tile; float32 fuses bias, residual and activation into the store, half writes float32 accumulators and an epilogue pass finishes. A float32 grid of fewer than 384 workgroups **splits k** (below) | | float32 on Metal, fast mode, windows and strided or small-image pointwise layers (channels per group a multiple of 16) | **taps-innermost implicit GEMM** (`ConvTap.hlsli`, 64 x 64 or, for small grids, 64 x 32 tiles), described under "Apple M2, float32, taps innermost" below; the pointwise layers of a normal image size keep the kernel of the next row | | float32 on Metal (simdgroup matrices) | **implicit GEMM on matrix hardware** (`ConvCoopF32.hlsl`): the 4-subgroup 64x64x16 simdgroup tile of `Gemm`, the window gathered while the input tile is staged (no column matrix), the bias, residual and activation applied as the result is stored, so the output is written once. A pointwise convolution (1x1, stride 1, no padding) whose tile lies inside the output reads weights and input straight from memory with no staging. Grids under 40 workgroups use a 64x32 tile. Deterministic mode, and fast mode where the row above does not apply (stem, depthwise-like channel counts, 1x1 layers of 64 or more pixels): both modes (see Determinism) | | float16 on Metal (simdgroup matrices) | the same implicit kernel with half operands (`ConvCoop.hlsli` with `CONV_F16`): float32 accumulation on one chain, bias, residual and activation applied as the halves are stored; with split k or an activation the kernel does not fuse, it leaves float32 planes in the accumulator scratch for the epilogue pass | | 1x1, stride 1, no padding, with matrix hardware (half) | `Gemm` on the input as it lies (no unfolding) | | any other shape with matrix hardware (half) | im2col (Chellapilla et al., 2006) into the column scratch, then `Gemm` (cooperative or simdgroup matrices) and the epilogue pass; groups loop | | transposed convolution | `Gemm` of the transposed weights with the input, then a gathering col2im (one thread per output element, no atomics) with the same epilogue | Variants are chosen once per instance from `GpuCapabilities.Compute`; a device without cooperative matrices (lavapipe, a forced-off run) runs the portable kernel, and the tests run every layer set in four configurations (fast, fast with matrix hardware off, deterministic, fast with the implicit kernel off). ### Not built, with the numbers - **Accumulation order.** The implicit-GEMM kernels sum k in blocks exactly as `Gemm` does (portable tile: 32-term block sums in 128-term groups; Apple matrix tile: a second accumulator over 128 terms); the im2col path inherits it from `Gemm`. Depthwise and direct kernels sum at most kernel-area terms in one chain. Scaled error against double at the Whisper stems and ResNet-50 3x3 layers is 1.4e-7 to 3.3e-7 (a single chain gave 4.6e-7 to 3.4e-6). - **Cooperative-matrix implicit GEMM** (window gather in the matrix tile load). On the RTX 4070, f16 ResNet-50 at b8 the im2col+`Gemm` layers already run at 6.6-13 TFLOP/s for 3x3 and 11-21 TFLOP/s for the 1x1 layers that use `Gemm` directly; the unfold pass is a bandwidth-bound write and read of at most a few tens of megabytes per layer. The ceiling for the change is the gap between the 3x3 and the 1x1 rate, about 1.3-1.6x on 3x3 layers, which are 2.4 of about 5.8 ms. Dropped as not worth a second coop kernel family now. - **Winograd F(2x2,3x3)** (Lavin and Gray, 2016); not needed once small grids split k (see above). `WinogradBound` in the tests times the sixteen batched products it needs with the real `Gemm` plus a copy moving the bytes its transforms must move: a floor no implementation can beat. RTX 4070, ResNet-50 stride-1 3x3 layers at b8 summed (host clock, indicative): **float32 direct 6.3 ms, floor 2.7 ms; half direct 2.4 ms, floor 1.0 ms.** The floor leaves room (the transforms run as three irregular passes, realistically 1.5x the copy floor, and half Winograd carries larger rounding error through the transformed data than the 2e-3 half tolerance leaves comfortably), but it is a speed-up for one model's 3x3 layers only, MobileNetV3 and DBNet have almost none. Left until ResNet-50 speed needs it; the variant would be a fourth `ConvVariant` behind the same `Encode`. ### Apple M2, float32 (indicative, host clock, 10 calls per command buffer; Metal has no timestamp queries) Convolution layers summed over the model's layer list, before and after the matrix-hardware implicit kernel (and the direct-load 4-subgroup `Gemm` it shares): | Model | batch | im2col + `Gemm` + epilogue pass | implicit simdgroup kernel | |---|---|---|---| | ResNet-50 | 1 | 26.0 ms | 24.4 ms | | ResNet-50 | 8 | 117.8 ms | 89.9 ms | | MobileNetV3-large | 1 | 5.8 ms | 4.3 ms | | DBNet 512 | 1 | 9.9 ms | 7.2 ms | The late layers of a classifier at batch 1 (14 x 14 and 7 x 7 images, 4608 deep) ran at 130-300 GFLOP/s: the grid had too few workgroups and each ran its 16-wide k steps one after another. They now **split k**: when the grid has fewer than 128 workgroups and k has at least 32 steps of 16, k is cut into up to 8 equal ranges (at least 16 steps each; the cut depends only on the shape). Each range writes a raw partial plane to the scratch, and the epilogue pass adds the planes in range order, then bias, residual and activation (so any activation works). No atomics. `Scratch()` reports the planes. Batch is part of the shape, so chunking an oversized batch does not change a bit. | Model (float32, fast, indicative) | batch | split-k off | split-k on | |---|---|---|---| | ResNet-50 | 1 | 24-26 ms | 20.3-22.3 ms (earlier session: 19 to 16.9) | | MobileNetV3-large | 1 | 4.24 ms | 3.50 ms | | DBNet 512 | 1 | no layer splits (k under 32 steps) | no layer splits; sums differ by run order on the shared GPU clock, not by kernel | | ResNet-50 | 8 | 90.0 ms | 90.4 ms | | MobileNetV3-large | 8 | 17.9 ms | 17.6 ms | The M2's clock state moves ResNet-50 b1 by about 20 % between sessions, so the table pairs runs made alternately in one session (five pairs; split-k won every one). Deterministic mode, ResNet-50 b1 20.4 ms, b8 90.5 ms. A version that prefetched the next tile into registers ran 40 % slower (register pressure), and wider narrow-tile thresholds ran slower at batch 8. ### RTX 4070, float32 portable kernel: split k for small grids At batch 1 the late layers of a classifier have few tiles (a 7 x 7 image is one tile of pixels, 512 output channels eight tiles) and each workgroup walks 4608 terms of k alone, so the 4070's 46 multiprocessors mostly idle: those layers ran at 0.4 to 1.5 TFLOP/s. The portable implicit kernel therefore splits k the way the Apple kernel does: when the grid has fewer than 384 workgroups and k has at least eight 32-term blocks, k is cut into up to 16 ranges (at least 128 terms each, aiming at about 512 workgroups) that run as workgroups of their own, each into its own plane of raw sums in the column scratch, and the epilogue pass adds the planes in range order, then bias, residual and activation (any activation works). Ranges are a multiple of 32 terms, so each starts on a block-sum boundary of the kernel; the cut depends only on the shape, so both modes use it and two runs are bit-identical. `Scratch()` reports the planes (`ranges * output size` per image; the static `ColumnElements` counts the unfolded columns only). A grid of 392 workgroups (56 x 56 at batch 8) stays whole: two ranges measured 4 to 30 % slower there. Used on every device whose float32 product is the portable tile: NVIDIA, AMD and Intel Vulkan, lavapipe, WebGPU (Metal has its own split above). Winograd, which this document priced at a floor of 2.3-2.4x, is not needed: ResNet-50's convolutions at batch 1 went from 6.48 to 2.57 ms. GPU time from timestamps, RTX 4070, driver 595.99, float32, fast mode, indicative (convolution layers summed over the model's layer list, each shape counted as often as the model uses it; whole-k is the same binary with the split turned off): | Model | batch | whole k | split k | |---|---|---|---| | ResNet-50 | 1 | 6.48 ms | 2.57 ms | | ResNet-50 | 8 | 16.67 ms | 15.52 ms | | MobileNetV3-large | 1 | 1.51 ms | 0.92 ms | | MobileNetV3-large | 8 | 2.87 ms | 2.57 ms | | DBNet 512 | 1 | 0.95 ms | 0.95 ms | | DBNet 512 | 8 | 6.77 ms | 5.75 ms | ResNet-50 batch 1, layers (microseconds, whole k then split k): 3x3 512 -> 512 at 7 x 7 520 -> 73; 3x3 512 -> 512 stride 2 on 14 x 14 532 -> 74; 3x3 256 -> 256 at 14 x 14 264 -> 70; 1x1 2048 -> 512 at 7 x 7 209 -> 36; 1x1 1024 -> 256 at 14 x 14 104 -> 36; 3x3 128 -> 128 at 28 x 28 138 -> 67. Scaled error against double over the 32 accumulation-test rows is at most 2.7e-7 (2.1e-7 to 3.3e-7 before). ### Apple M2, float32, taps innermost (indicative, GPU timestamps, 30 calls per command buffer) `ConvCoop.hlsli` numbers k as (input channel, kernel row, kernel column), so each staged element needs a division chain, two bounds tests and a tap update, and the float32 wide tile staged without prefetch. Measured with a register-only kernel the M2's simdgroup matrices do 3.25 TFLOP/s in float32 (90 % of the ALU peak) and a plain 64 x 64 matrix product with the blocked sums above does 1.9 TFLOP/s, so the convolutions' 0.4-0.8 TFLOP/s were staging cost, not matrix throughput. `ConvTap.hlsli` runs k in chunks of 16 input channels and, inside a chunk, over the taps of the window: a step of 16 is one tap of 16 channels, so the tap's row, column and padding test are the same for every element a thread stages, the channels' addresses differ by one input plane, and a thread's step is one test and eight independent loads (plus eight weight loads, next step in flight). Bias, residual and activation are still applied as the result is stored. It is used for float32 in fast mode when the channels per group are a multiple of 16, for every layer except pointwise ones with at least 64 (32 for the narrow tile) pixels, whose direct fragment loads measured at 1.4-1.5x the speed of staging. Deterministic mode and every other shape keep the previous kernel, so deterministic results are unchanged bit for bit (the two-run hashes of the conv tests equal those of the previous build); the fast sums differ in order only, with the same block structure. Convolution layers of ResNet-50 summed by class (ms; each shape counted as often as the model uses it; the stem, 7 x 7 with 3 input channels, is untouched): | Class | b1 before | b1 now | b8 before | b8 now | |---|---|---|---|---| | 3x3 stride 1 | 5.08 | 4.50 | 45.2 | 27.2 | | 3x3 stride 2 | 1.14 | 1.06 | 11.5 | 6.4 | | 1x1 stride 2 (downsample) | 1.42 | 0.90 | 11.7 | 6.3 | | 1x1 stride 1 (reduce, expand) | 4.36 | 4.57 | 32.3 | 27.9 | | stem | 0.50 | 0.50 | 3.7 | 3.7 | | **sum** | **12.4** | **11.5** | **104** | **71.5** | Scaled error against float64 on the 23 ResNet-50 layers at batch 1 and 8, the MobileNetV3 and DBNet layers and the 35 edge shapes: at most 5.9e-7 (3.8e-7 for ResNet-50), the same as the previous kernel. Whole model through onemodel's `.Gpu` (fast mode, 10 runs, same session, alternating builds): ResNet-50 float32 b1 14.55 -> 13.15 ms, b8 110.9 -> 78.1 ms, scaled error against onemodel's CPU exact mode 1.28e-6 (before: 1.3e-6). What is left between this and a Core ML run of the same model on the same GPU (6.1 ms at b1, 35.3 ms at b8; its float32 output differs from ours by 1e-6 as ours does): - Ours is 0.9 TFLOP/s over the convolutions at b8 (0.7 at b1), against 1.9 TFLOP/s for a plain matrix product, and a run at 35 ms means 1.9 TFLOP/s over the whole 65 GFLOP model: more than a direct convolution can reach on this GPU, so the reference must also save arithmetic (Winograd or similar) on the 3 x 3 layers (45 % of the model's flops). - The weights of a 3 x 3 layer are read as 16 words at a stride of 9 per row and step; with the weights re-laid out so those reads are contiguous (a timing-only change, results wrong) the 3 x 3 layers ran at 1.27x the speed. A real layout needs a per-model weight preparation pass, which this API does not have. - Table-driven window decode with the weights in their own order (contiguous weight reads) cost more per staged element than it saved: b8 71.5 -> 89.9 ms. - At batch 1 the layers are latency bound (12 to 140 workgroups, 4 to 150 steps each): moving the narrow-tile and split-k thresholds in either direction measured slower (best 11.8 against 12.6-14.0 ms). - Pointwise layers with 64 or more pixels still use the direct-load kernel at 0.6-1.6 TFLOP/s; those with few input channels are bound by the epilogue (tile through workgroup memory, scalar stores). ## Pooling and resize - `MaxPool` / `AveragePool`: padding, dilation, `ceil_mode` (a last window that starts inside the end padding is dropped, as in ONNX), `count_include_pad` (padding counts, taps beyond it do not). Max skips NaN taps unless all are NaN; an empty max window gives the lowest finite value of the type (-3.4028235e38; -65504 for half); an empty average window gives NaN. `MaxPool`'s indices output is not provided. `GlobalAveragePool` / `GlobalMaxPool` reduce a plane with a fixed halving tree. - `Resize`: nearest (all four rounding modes) and linear along each axis (bilinear in two), coordinate transformations half_pixel, half_pixel_symmetric, pytorch_half_pixel, align_corners and asymmetric. The float32 coordinate arithmetic runs on the host once per shape into tap tables (`ResizePlan`, create once, encode many times) so every device gets the same taps. Not provided: cubic, antialias, `exclude_outside`, `tf_crop_and_resize`. - `BatchNormalization` has no kernel: inference batch norm folds into the convolution weights in the graph, and the DBNet and reference exports are already folded. ## ConvInteger `ConvIntegerKernels` is ONNX `ConvInteger`: `y = sum (x - xZero) * (w - wZero[m])` in int32 with wrap-around, for uint8 and int8 `x` and `w` in any mix, the geometry of `ConvProblem` (strides, dilations, asymmetric padding, groups, 1-D as height 1). The zero point of `x` is one value and that of `w` one value or one per output channel, in device buffers of their operand's own type. A padded tap is the real value zero, which is the quantized value `xZero`, so it adds nothing: the unfolding pass writes `xZero` into padded positions and the product subtracts it again, which is how onemodel's CPU kernel treats padding. It is lowered to `Int8Gemm` as the float convolution is to `Gemm` (Chellapilla et al., 2006): one pass unfolds the windows of each image and group into a byte matrix (`K = InChannels / Groups * KH * KW` rows, rows padded to a multiple of four bytes so the product reads whole words), then the weights of the group multiply it. One group makes one product for the whole batch. Several groups make one product per image and group, which is slow for depthwise convolutions with many channels (no direct 8-bit depthwise kernel yet). Sums are exact integers, so the packed-dot and portable products, every device and both `GpuKernelMode` values give the same bits. Column scratch is the instance's own (up to 64 MiB, then the batch runs in chunks of images) or caller-placed (`GpuBuffer` and a byte offset that is a multiple of four, `ConvIntegerKernels.ScratchBytesPerImage`). Tested bit for bit against a plain C# definition of the operator on the Whisper-small front end (80 to 768 channels over 3000 frames, kernel 3, and 768 to 768 with stride 2), ResNet-50 and MobileNet shapes, dilation, asymmetric padding, groups, offsets, extreme values and all four type combinations, with the packed-dot product and with it ruled out. ## Tolerances (what the tests assert; scaled error as in [kernels.md](https://recursiveai.org/onegpu/docs/kernels/)) | Operation | float32 | half in, float32 accumulation | |---|---|---| | Conv, ConvTranspose (fast and deterministic) | 1e-5 (measured <= 2.1e-6 over 23 ResNet-50, 45 MobileNetV3-large and 34 DBNet layer shapes, b1 and b8) | 2e-3 (measured <= 4.4e-4) | | ConvInteger | bit-exact (int32) | | | MaxPool, GlobalMaxPool | bit-exact | bit-exact | | AveragePool, GlobalAveragePool | 1e-5 (measured 2.3e-7) | 2e-3 | | Resize nearest | bit-exact | bit-exact | | Resize linear | 1e-5 (measured 9.8e-8) | 2e-3 (measured 2.8e-4) | ## Determinism No kernel here uses an atomic or a subgroup operation. Each output is summed in an order fixed by the shape: the matrix kernel's, the window order of the direct kernels, or, for split-k, shape-determined ranges of k added in range order. On Apple GPUs `GpuKernelMode.Deterministic` runs the float32 simdgroup implicit kernel too (a simdgroup matrix product is a fixed function of its operands, and each output tile is summed by one simdgroup in k order), which runs at 1.6-1.7x the speed of the scalar tile it used before (ResNet-50 b1 27-29 ms to 16.4-17.1 ms, b8 145 to 90 ms, MobileNetV3 b1 5.2 to 3.8 ms); elsewhere it selects the portable matrix kernel. Results are bit-identical run to run on one device and driver in that mode; `Fast` may use matrix hardware on other devices, whose summation order is the hardware's. The tests run two passes of seven layer shapes (two split) per element type and compare hashes. ### Apple M2, float16 and the pointwise edge tile (indicative, GPU timestamps, 100 calls per command buffer) The last pixel tile of a pointwise layer whose pixel count is not a multiple of 64 (14 x 14 = 196, 28 x 28 = 784) is moved back to end at the last pixel, so it is a whole tile and reads its operands directly. It recomputes pixels the previous tile also writes, from the same inputs in the same order, so the stores agree. Float32 with the wide tile stages straight into the tiles (register staging cost occupancy there: ResNet-50 b8 127 ms against 110); the narrow tile keeps it. | Convolution layers of | batch | float32 before | float32 now | float16 before | float16 now | |---|---|---|---|---|---| | ResNet-50 | 1 | 17.0 ms | 12.4 ms | 18.0 ms | 9.7 ms | | ResNet-50 | 8 | 110 ms | 105 ms | 114 ms | 62 ms | Where the float16 time goes at batch 1 (9.7 ms): 3x3 layers 4.8 ms at about 800 GFLOP/s, 1x1 stride-1 layers about 3.5 ms (600-1400), stride-2 1x1 projections about 1 ms, the stem 0.3 ms; the measured matrix peak of the M2 is about 1.7 TFLOP/s. The 3x3 layers are bound by the gathering of the input tile (one scattered load per element); loading the weights straight from memory (no change) and compile-time 3 x 3 window arithmetic (9% on those layers) do not close it. A kernel that reads padded input rows as matrix operands would need a padded copy of the input and a weight layout with the taps outermost. # Recurrent layers `RecurrentKernels` runs ONNX `RNN`, `GRU` and `LSTM` layers on the GPU: forward, reverse and bidirectional, layouts 0 and 1, `sequence_lens`, initial states, peepholes, `clip`, `input_forget`, `linear_before_reset`, and every activation ONNX names (Sigmoid, Tanh, Relu, Affine, LeakyRelu, ThresholdedRelu, ScaledTanh, HardSigmoid, Elu, Softsign, Softplus, with `activation_alpha` and `activation_beta` already resolved by the caller into `GpuRecurrentActivation`). Tensors are float32 or `Half` (all of them the same type) with float32 accumulation and float32 state. The semantics are those of onemodel's CPU kernels, which settled what the specification leaves open. ```csharp using var rnn = new RecurrentKernels(device); // GpuKernelMode.Fast by default var problem = new RecurrentProblem(GpuRecurrentCell.Lstm, seq: 80, batch: 1, input: 256, hidden: 256) { Direction = GpuRecurrentDirection.Bidirectional }; var tensors = new RecurrentOperands { X = new(x), W = new(w), R = new(r), B = new(b), Y = new(y), YH = new(yh) }; rnn.Encode(encoder, problem, tensors); ``` W, R, B, P are used as stored (`[dirs, gates * hidden, input]` and so on, gate order i, o, f, c for LSTM and z, r, h for GRU); there is no repacking. Every tensor is a buffer plus an element offset (`GpuTensorRef`). Outputs that are null are not written. `Half` tensors follow the usual packing rule (buffers end on a 32-bit word, `GpuKernelBuffers.PaddedHalfLength`); an output that shares a word with a neighbour keeps the neighbour's half. ## How a layer runs 1. **Input projection** of every step and both directions: one batched `Gemm` (`Xw = X · Wᵀ`, `Batch` = directions, W through the transposed-B form). The family creates the `Gemm` for its mode, so `Fast` may use matrix hardware for `Half` and `Deterministic` uses the portable tile. 2. **State setup**: one pass that writes the first hidden state and the cell state from `InitialH`/`InitialC` or zeros. 3. **One pass per step** (two for a GRU with `linear_before_reset = 0`, because the candidate's product needs `r ⊙ H` of every unit). A pass covers both directions and all batch rows. A workgroup owns 32 rows of R (hidden units × gates of the pass) and, for each, 8 lanes sum `k = lane, lane + 32, …` in increasing order; the 32 partial sums are combined by a halving tree in shared memory; then one thread per (unit, batch row) does the gate arithmetic. The hidden state is two buffers that swap every step, so no workgroup reads what another writes in the same pass; the barrier between passes orders the steps. Rows whose sequence has ended copy their state through. The thread that forms a unit's gates adds the 8 lane sums of each row itself, in the order of the halving tree (`((p0 + p4) + (p2 + p6)) + ((p1 + p5) + (p3 + p7))`), after one barrier. The kernel is compiled per cell (LSTM, GRU, RNN), per precision, for one batch row or up to eight at a time, and for the ONNX default activations (Sigmoid, Tanh) or all eleven. 4. **One pass per requested output** (Y, Y_h, Y_c): layout, zeroing past the sequence length, rounding to `Half`. Why a pass per step: a kernel that loops over all steps in one dispatch needs every workgroup of the dispatch to reach a grid-wide barrier, which Vulkan, Metal and WebGPU do not guarantee (workgroups may not all be resident). The alternative that needs no grid barrier, one workgroup per direction that walks every step itself ordered by workgroup barriers, was built and measured and is not shipped: a workgroup runs on one compute unit, so each step's 1 MB of weights comes through one unit's memory path. Bidirectional LSTM, sequence 80, batch 1, input and hidden 256, float32, untimed, same bits as the pass per step: Apple M2 9.9 ms against 2.4 ms for a pass per step, RTX 4070 5.1 ms against 0.7 ms. What a step costs is mostly the kernel's start, not its arithmetic. A step kernel that held all eleven activation functions inline (five call sites, a switch each) cost about 20 µs more per step on an Apple M2 than the same kernel without them (removing the activations alone took 1.5 ms off an 80-step layer; the size of the kernel, not its arithmetic, is the likely cause); the default activations get a kernel without them, and each cell gets a kernel without the other cells. Carrying eight batch rows through the product when the batch is one cost another 2x (registers and shared memory limit how many workgroups are resident); a batch of one gets a one-row kernel. `RecurrentProblem.MaxLength` (a host-side bound on the sequence lengths) shortens the loop when `sequence_lens` is given. Scratch: the input projection (`RecurrentKernels.XwLength` floats) and the state with every step's output (`StateLength` floats) are kept by the family, grown on demand, or supplied by the caller (`xwScratch`, `stateScratch`, each with an element offset `xwScratchOffset`, `stateScratchOffset`, so both regions may sit in one arena buffer; a browser refuses one buffer bound read-only and read-write in a dispatch, so there they stay in two buffers) so a memory planner can place them. Both must fit one storage binding; larger layers throw `GpuUnsupportedException`. Sizes, offsets and shapes are checked in safe arithmetic before anything is recorded (`GpuKernelShapeException` naming `LSTM`/`GRU`/`RNN` and the argument). ## Arithmetic and determinism The order is fixed by the shapes: the projection is `Gemm`'s; each recurrent dot product is the 32-lane strided sums then the halving tree; the gate arithmetic is ONNX's `pre = ((xw + rec) + Wb) + Rb`, then the peephole terms (`i`, `f` with the previous cell state, `o` with the new one), then `clip`, then the activation, then `c = f·c + i·c̃`, `h = o·h_act(c)` (GRU: `h = (1 − z)·h̃ + z·h`). No atomics, no subgroup arithmetic, nothing depends on the subgroup size or scheduling. One kernel therefore serves both `GpuKernelMode` values and a run is bit-identical to the next on one device and driver (checked: every test case runs twice; the recognizer shapes also in `Deterministic` mode). Whether a compiler fuses a multiply with the following add is the device's, so results differ in the last bits between devices; they are compared with the reference below, not with each other. ## Semantics - `sequence_lens`: row `b` runs its first `len` steps, in reverse from step `len − 1`; Y is zero at the steps past `len`; Y_h and Y_c are the state after the last valid step; **a row of length 0 has Y_h = Y_c = 0** (onemodel's CPU kernel; the specification is silent). Lengths are int32 on the device and not validated there (a value above `Seq` is treated as `Seq`); validate them on the host. - `Seq = 0`: nothing runs; Y_h and Y_c are the initial states or zeros. `Batch = 0`: nothing is recorded. - `clip` clamps every gate's pre-activation to `[−clip, clip]` when positive (NaN passes through). - `input_forget` couples the forget gate: `f = 1 − i`. `linear_before_reset = 1`: `h̃ = g(xw_h + r·(rec_h + Rb_h) + Wb_h)`. - Peepholes, `input_forget`, `InitialC` and `YC` exist only for an LSTM, `linear_before_reset` only for a GRU; anything else throws. ## Tolerance and evidence Scaled error `max|gpu − ref| / max|ref|` against a double-precision C# reference written from the operator definitions: float32 ≤ 1e-5, `Half` (float32 accumulation and state, outputs rounded to half) ≤ 2e-3. The tests cover every cell × direction × layout with option sets (bias, lengths including 0, initial state, peepholes, clip, `input_forget`, `linear_before_reset`), every activation in every slot of both directions, odd offsets with sentinels around every output, empty sequence and batch, hidden sizes that do not divide the workgroup, batches above eight, and the shapes of a text-line recognizer (bidirectional LSTM, sequence 80 and 48, batch 1 and 8, input 256 and 192, hidden 256 and 96). # Speed at 1.0 > **Note:** in the README’s words, these figures were “Measured through onemodel’s `RecursiveAI.OneModel.Gpu`”. That package isn’t released yet; OneModel’s GPU support is coming in a later OneModel release, so the figures can’t be reproduced with the released packages. The harness and the full method are in onemodel-csharp, `tools/bar.sh` and `docs/bench.md`. The models are onemodel's reference set: MiniLM-L6, BERT-base, ResNet-50, MobileNetV3, DBNet, CRNN and Whisper-small (encoder and decoder), in float32, float16 and dynamically quantized 8-bit forms. Speed is judged against the reference provider on the same machine and the same model file; ratio is ours divided by the reference, so lower is better. The bar is 1.5 on a desktop (target 1.15) and 2.5 in a browser (target 1.5). `met` is a ratio above the target and within the bar, `target` is a ratio within the target, `missed` is above the bar. Only **gate rows** are judged (the reference ran every compute node on its GPU, read from its own placement log and, for Core ML, its compute plan); **information rows** are listed with their ratio and never judged. Every row runs a model through onemodel's `RecursiveAI.OneModel.Gpu` on these packages at their 1.0.0 code, fast mode, validation layers off, host arrays in and out, the median of warm runs (Linux 20, Apple M2 10, Chrome 10 per page) in one benchmark-queue job per group of rows, each group alone on its machine. Load average sampled every 5 s before, during and after is in every row of the jsonl; the one-minute load average at the start of a group still carries the harness's own CPU reference runs of the group before it. Reference versions: ONNX Runtime 1.31.1 with the `onnxruntime-ep-webgpu` 0.4.0 plugin (Dawn over Vulkan on Linux, over Metal on macOS), Core ML through the same wheel on GPU compute units, onnxruntime-web 1.31.1. ### Linux, RTX 4070 (Vulkan, driver 595.99.02), Ryzen 9 7950X3D, 32 threads; reference: WebGPU provider Three groups of rows, each alone on the machine. Gate rows 14 of 20: **all met**, worst ResNet-50 float32 at batch 8 (1.40) and BERT float32 at batch 8 (1.35); the other twelve are at or below the 1.15 target. The int8 rows are information rows: the reference leaves MatMulInteger, DynamicQuantizeLinear, QuantizeLinear or ConvInteger to its CPU provider, so those rows are not a like-for-like comparison. CUDA and onemodel CPU columns were not run for the gate. | row | ours ms (GPU kernels) | reference | ref ms | ours / ref | class | verdict | other provider | |---|---|---|---|---|---|---|---| | minilm_fp32_b1 | 1.69 (1.17) | WebGPU EP | 4.309 | 0.39 | gate | target | - | | minilm_fp32_b8 | 5.5 (4.83) | WebGPU EP | 6.4 | 0.86 | gate | target | - | | minilm_fp16_b1 | 2.24 (1.42) | WebGPU EP | 3.982 | 0.56 | gate | target | - | | minilm_int8_b1 | 2.09 (1.37) | WebGPU EP | 48.566 | 0.04 | information | target | - | | bert_fp32_b1 | 6.49 (5.58) | WebGPU EP | 7.471 | 0.87 | gate | target | - | | bert_fp32_b8 | 31.19 (30.35) | WebGPU EP | 23.188 | 1.35 | gate | met | - | | bert_int8_b1 | 6.33 (5.05) | WebGPU EP | 115.786 | 0.05 | information | target | - | | resnet50_fp32_b1 | 3.18 (2.7) | WebGPU EP | 3.73 | 0.85 | gate | target | - | | resnet50_fp32_b8 | 16.44 (15.21) | WebGPU EP | 11.773 | 1.40 | gate | met | - | | resnet50_fp16_b1 | 3.87 (3.07) | WebGPU EP | 4.019 | 0.96 | gate | target | - | | resnet50_int8_b1 | 4.36 (3.39) | WebGPU EP | 104.208 | 0.04 | information | target | - | | mobilenet_fp32_b1 | 1.45 (0.99) | WebGPU EP | 1.899 | 0.76 | gate | target | - | | mobilenet_fp32_b8 | 3.53 (2.74) | WebGPU EP | 3.666 | 0.96 | gate | target | - | | mobilenet_int8_b1 | 3.23 (1.96) | WebGPU EP | 112.891 | 0.03 | information | target | - | | dbnet_fp32_b1 | 2.48 (1.8) | WebGPU EP | 3.6 | 0.69 | gate | target | - | | crnn_fp32_b1 | 0.95 (0.52) | WebGPU EP | 19.594 | 0.05 | gate | target | - | | whisper_enc_b1 | 57.4 (56.1) | WebGPU EP | 52.142 | 1.10 | gate | target | - | | whisper_dec_b1 | 14.37 (12.38) | WebGPU EP | 13.55 | 1.06 | gate | target | - | | whisper_enc_int8_b1 | 32.88 (31.46) | WebGPU EP | 650.538 | 0.05 | information | target | - | | whisper_dec_int8_b1 | 10.45 (7.95) | WebGPU EP | 184.255 | 0.06 | information | target | - | ### Apple M2 (Metal, macOS 27.0.1); reference: whichever gate configuration of Core ML (GPU compute units) and the WebGPU provider on Metal took less time Three groups in the benchmark queue, 10 runs per row. Gate rows 14 of 20: twelve are met, **two are misses for 1.0: ResNet-50 float32 at batch 1 (2.16) and batch 8 (2.21) against Core ML on the GPU** (6.07 ms and 35.27 ms; Core ML MLProgram on the GPU computes in float32 to 1e-6, so it is a fair reference). Against the WebGPU provider on Metal the same rows are 0.82 (13.09 against 15.94 ms) and 0.81 (77.94 against 95.71 ms). The miss is documented and ships: the levers left are per-model weight preparation for contiguous 3x3 reads and a Winograd F(2x2,3x3) kernel, estimated at about 1.8 times Core ML at batch 8, so not enough to close it. MobileNetV3 float32 at batch 8 (1.16) is just above the target. The int8 rows are information rows (MatMulInteger, DynamicQuantizeLinear and ConvInteger on the reference's CPU provider). | row | ours ms (GPU kernels) | reference | ref ms | ours / ref | class | verdict | other provider | |---|---|---|---|---|---|---|---| | minilm_fp32_b1 | 6.67 (6.19) | WebGPU EP | 9.769 | 0.68 | gate | target | Core ML (GPU) 10.404 (information) | | minilm_fp32_b8 | 25.91 (25.13) | WebGPU EP | 40.709 | 0.64 | gate | target | Core ML (GPU) 72.514 (information) | | minilm_fp16_b1 | 10.61 (9.72) | WebGPU EP | 9.55 | 1.11 | gate | target | Core ML (GPU) 11.91 (information) | | minilm_int8_b1 | 12.65 (11.61) | Core ML (GPU) | 6.917 | 1.83 | information | missed | WebGPU EP 18.781 (information) | | bert_fp32_b1 | 24.7 (23.99) | WebGPU EP | 41.456 | 0.60 | gate | target | Core ML (GPU) 66.332 (information) | | bert_fp32_b8 | 156.11 (155.02) | WebGPU EP | 221.278 | 0.71 | gate | target | Core ML (GPU) 512.805 (information) | | bert_int8_b1 | 49.59 (48.7) | Core ML (GPU) | 29.128 | 1.70 | information | missed | WebGPU EP 53.508 (information) | | resnet50_fp32_b1 | 13.09 (12.5) | Core ML (GPU) | 6.068 | 2.16 | gate | missed | WebGPU EP 15.937 (gate) | | resnet50_fp32_b8 | 77.94 (76.77) | Core ML (GPU) | 35.268 | 2.21 | gate | missed | WebGPU EP 95.712 (gate) | | resnet50_fp16_b1 | 13.13 (12.41) | WebGPU EP | 12.462 | 1.05 | gate | target | Core ML (GPU) 5.409 (information) | | resnet50_int8_b1 | 21.87 (20.79) | Core ML (GPU) | 18.541 | 1.18 | information | met | WebGPU EP 69.618 (information) | | mobilenet_fp32_b1 | 5.33 (4.83) | WebGPU EP | 5.045 | 1.06 | gate | target | Core ML (GPU) 9.886 (information) | | mobilenet_fp32_b8 | 21.04 (19.88) | WebGPU EP | 18.072 | 1.16 | gate | met | Core ML (GPU) 28.363 (information) | | mobilenet_int8_b1 | 11.17 (9.89) | Core ML (GPU) | 13.924 | 0.80 | information | target | WebGPU EP 50.786 (information) | | dbnet_fp32_b1 | 17.47 (16.44) | WebGPU EP | 17.859 | 0.98 | gate | target | Core ML (GPU) 13.791 (information) | | crnn_fp32_b1 | 2.79 (2.37) | WebGPU EP | 48.142 | 0.06 | gate | target | Core ML (GPU) 1.534 (information) | | whisper_enc_b1 | 371.95 (370.23) | WebGPU EP | 476.461 | 0.78 | gate | target | Core ML (GPU) 1047.15 (information) | | whisper_dec_b1 | 88.27 (85.95) | WebGPU EP | 79.586 | 1.11 | gate | target | Core ML (GPU) 132.07 (information) | | whisper_enc_int8_b1 | 609.78 (607.89) | WebGPU EP | 563.7 | 1.08 | information | target | Core ML (GPU) 567.703 (information) | | whisper_dec_int8_b1 | 86.91 (85.48) | Core ML (GPU) | 48.924 | 1.78 | information | missed | WebGPU EP 104.941 (information) | ### Chrome 155 on the RTX 4070, WebGPU (ahead-of-time compiled page); reference: onnxruntime-web WebGPU provider Two groups, each alone on the machine, ten warm runs per page, same adapter. The twelve float32 rows; every one is within the 2.5 bar and eleven within the 1.5 target (ResNet-50 batch 8, 1.54, is the only one above). The float16 rows decline in this browser (no `shader-f16` on this Chrome) and the int8 rows are information, so none is a bar row here; they run (fast and exact, two-run output identical) in the device matrix below. | row | ours ms | onnxruntime-web ms | ours / ref | verdict (bar 2.5, target 1.5) | |---|---|---|---|---| | minilm_fp32_b1 | 3.8 | 7.2 | 0.53 | target | | minilm_fp32_b8 | 10 | 9.3 | 1.08 | target | | bert_fp32_b1 | 10.6 | 10.4 | 1.02 | target | | bert_fp32_b8 | 38.6 | 26.8 | 1.44 | target | | resnet50_fp32_b1 | 4.8 | 6.2 | 0.77 | target | | resnet50_fp32_b8 | 21.6 | 14 | 1.54 | met | | mobilenet_fp32_b1 | 3.7 | 5.8 | 0.64 | target | | mobilenet_fp32_b8 | 7.5 | 9 | 0.83 | target | | dbnet_fp32_b1 | 5.2 | 9.3 | 0.56 | target | | crnn_fp32_b1 | 3.2 | 21.2 | 0.15 | target | | whisper_enc_b1 | 68.4 | 56.6 | 1.21 | target | | whisper_dec_b1 | 22.3 | 17.7 | 1.26 | target | ### Windows, Intel Iris Xe (Windows 11, driver 101.7085), indicative only Not publishable: one run on a shared laptop, ONNX Runtime 1.24.4 DirectML wheel, `tools/bar.ps1`. The float32 transformer rows (BERT, MiniLM), ResNet-50 (float32 and float16), MobileNetV3 at batch 8 and Whisper **miss the 1.5 bar against DirectML (1.53 to 2.89)**: that GPU has no cooperative matrix, so the portable tiled kernels run. They ship as documented misses. NativeAOT on Windows is untested. | row | ours ms | DirectML ms | ours / DirectML | WebGPU EP ms | ours / WebGPU EP | |---|---|---|---|---|---| | bert_fp32_b1 | 116.22 | 42.543 | 2.73 | 132.137 | 0.88 | | bert_fp32_b8 | 871.4 | 301.285 | 2.89 | 577.978 | 1.51 | | bert_int8_b1 | 47.61 | 129.402 | 0.37 | 281.994 | 0.17 | | crnn_fp32_b1 | 5.41 | 18.992 | 0.28 | 81.591 | 0.07 | | dbnet_fp32_b1 | 34.32 | 27.237 | 1.26 | 41.974 | 0.82 | | minilm_fp16_b1 | 21.06 | 22.841 | 0.92 | 16.371 | 1.29 | | minilm_fp32_b1 | 28.64 | 10.917 | 2.62 | 19.652 | 1.46 | | minilm_fp32_b8 | 113.63 | 45.324 | 2.51 | 92.67 | 1.23 | | minilm_int8_b1 | 15.7 | 23.65 | 0.66 | 34.099 | 0.46 | | mobilenet_fp32_b1 | 11.72 | 9.625 | 1.22 | 11.516 | 1.02 | | mobilenet_fp32_b8 | 50.23 | 32.534 | 1.54 | 42.67 | 1.18 | | mobilenet_int8_b1 | 16.78 | 48.483 | 0.35 | 481.871 | 0.03 | | resnet50_fp16_b1 | 79.69 | 30.107 | 2.65 | 30.548 | 2.61 | | resnet50_fp32_b1 | 51.16 | 33.364 | 1.53 | 40.245 | 1.27 | | resnet50_fp32_b8 | 375.45 | 141.546 | 2.65 | 216.923 | 1.73 | | resnet50_int8_b1 | 63.39 | 88.177 | 0.72 | 539.153 | 0.12 | | whisper_dec_b1 | 315.7 | 139.074 | 2.27 | 203.836 | 1.55 | | whisper_dec_int8_b1 | 124.29 | 312.933 | 0.4 | 1135.404 | 0.11 | | whisper_enc_b1 | 1843.84 | 789.523 | 2.34 | 1388.869 | 1.33 | | whisper_enc_int8_b1 | 960.42 | 1051.38 | 0.91 | 5388.266 | 0.18 | ### Correctness of the same runs `OneModel.Gpu.Tests` (383 tests, 0 failed, 0 skipped): RTX 4070 and llvmpipe and RADV (Mesa 25.2.8) with the Khronos validation layer on, 0 validation messages; Apple M2 with Metal API and shader validation fatal. All 21 rows (20 above plus the greedy decode) with `RequireGpu`, fast and exact: 100 % of nodes in one region on every device, outputs identical between two runs in each mode (40 of 40 model-mode pairs per device), 0 validation messages. float32 rows are within a scaled error of 1.03e-5 of the CPU's exact mode on every device (Whisper-small encoder 8.0e-5), float16 within 1.2e-3, the 8-bit rows within the int8 policy (model-level error against the CPU exact mode 3.9e-4 to 8.8e-1, reported next to the CPU's own fast-versus-exact drift). CRNN greedy ids and the Whisper decoder's greedy tokens (8 of 8) equal the CPU's on every device. The Whisper-small encoder at batch 4 against a float64 evaluation, 16 inputs: mean error 1.14 (fast) and 1.21 (exact) times the CPU exact mode's on the RTX 4070 (max 0.90 and 0.97 times), 1.11 and 1.22 on the M2 (max 1.17 and 1.00), within 1.5 times the CPU exact mode's error. Chrome (RTX 4070): all rows fast and exact with `RequireGpu`, one region, identical runs, no page errors; SwiftShader subset (MiniLM, DBNet, CRNN, ResNet-50, Whisper decoder) likewise. `OneModel.Tests` 999 passed, 4 skipped, golden table equal to the base, NativeAOT smoke publish without trim or AOT warnings, run on RTX 4070 and llvmpipe with validation (all ConvInteger and model checks `ok`). # Changelog All notable changes are listed here. The format follows [Keep a Changelog](https://keepachangelog.com/en/1.1.0/) and the project follows [Semantic Versioning](https://semver.org/). ## [1.0.0] - 2026-11-12 First stable release. Everything below is new in 1.0; the public API of every package is recorded in its `PublicAPI.Shipped.txt`, and names (package ids, namespaces, `ONEGPU_*` variables, the `onegpu` JavaScript module) are listed in [`docs/environment.md`](https://recursiveai.org/onegpu/docs/environment/). ### Device layer - `RecursiveAI.OneGpu`: `GpuDevice`, adapters and `GpuCapabilities` as data (subgroup sizes and operations, float16, bfloat16, int8 and integer dot product, cooperative-matrix shapes, device limits), typed `GpuBuffer`, uploads, non-blocking `ReadAsync` readback, compute pipelines checked against the shader's reflection, command encoders with dispatch, indirect dispatch, pass timestamps and `ComputePassCount` / `DispatchCount`, memory statistics, typed errors that name the capability, what was expected and what the device has. Surfaces, textures, graphics pipelines and render passes for renderers. - Backends: Vulkan 1.3 (own memory allocator, timeline-semaphore readback, validation routing per device), Metal on Apple silicon (typed `objc_msgSend`, shared compute encoders), WebGPU in the browser (`[JSImport]`, several devices per page). `GpuBackends` picks one; `ONEGPU_BACKEND` overrides. - No native binaries of ours. The core and the kernels have no interop. All packages are trimmable and Native AOT compatible. ### Kernels - GEMM in float32, float16, bfloat16 (cooperative-matrix variants on Vulkan and Metal where the device has them) and int8, with fused bias and activation epilogues; fused attention with masks, causal rule, grouped heads and Apple simdgroup variants; Conv, ConvTranspose, ConvInteger, pooling and resize; RNN, GRU and LSTM; broadcasting element-wise, integer and bool operations; reductions, softmax, layer and RMS normalisation; gather, scatter, slice, concat, expand and transpose; quantize, dequantize and dynamic quantize. - Float32 sums accumulate in blocks, so their error against a double-precision sum is roughly ten times lower than a single chain's. - `GpuKernelMode.Deterministic` for GEMM, reductions, softmax and attention: a fixed order, no float atomics, bit-identical runs on one GPU and driver. Tolerances per operation and dtype are in [`docs/kernels.md`](https://recursiveai.org/onegpu/docs/kernels/). - Shaders are HLSL compiled at build time for SPIR-V, Metal Shading Language and WGSL, so a layout mismatch is a build error. ### Tooling - Tests run with the Khronos validation layer required and clean on Vulkan, with API and shader validation fatal on Metal, and in headless Chrome (SwiftShader and a hardware adapter) on WebGPU. `tools/ci` holds the scripts. - `PublicAPI.Shipped.txt` per package makes any public change a build error until the file changes in the same diff. [1.0.0]: https://github.com/recursive-ai-oss/onegpu-csharp/releases/tag/v1.0.0