# oneshader docs, all pages (OneShader 1.2.0) Every page of https://recursiveai.org/oneshader/docs/ and https://recursiveai.org/oneshader/cookbook/ as markdown, getting started first. Index: https://recursiveai.org/oneshader/llms.txt # Getting started with OneShader Compile HLSL shaders during `dotnet build` with errors on the HLSL line, use the C# structs and bindings generated from them, emit Metal and WGSL next to SPIR-V, build pipeline layouts, and compile at run time or from the command line. OneShader is an HLSL compiler written in C# with no dependency beyond .NET itself. Modern HLSL (Shader Model 6 style, HLSL 2021) for vertex, pixel and compute shaders goes in. SPIR-V for Vulkan 1.2 and newer, Metal Shading Language and WGSL come out, each with reflection. Every sample on this page was compiled and run against the 1.2.0 packages from NuGet. ## Install OneShader targets .NET 10 (`net10.0`), so you need the .NET 10 SDK. For shaders built into a project, reference the build package: ```sh dotnet new console -o Start cd Start dotnet add package RecursiveAI.OneShader.Build --version 1.2.0 ``` Build with the .NET SDK (`dotnet build` or `dotnet run`). Visual Studio’s .NET Framework MSBuild can’t load the build task. The build package brings the compiler package, `RecursiveAI.OneShader`, with it. The license is MIT or Apache-2.0, at your choice. ## Compile a shader with the project Every `*.hlsl` file in the project is compiled on build. Save this as `Shaders/Sprite.hlsl`: ```hlsl cbuffer Frame : register(b0) { float4x4 ViewProjection; float4 Tint; }; [[vk::binding(1)]] Texture2D Atlas; [[vk::binding(2)]] SamplerState AtlasSampler; struct VSIn { float2 Position : POSITION; float2 Uv : TEXCOORD0; [[oneshader::format("unorm8x4")]] float4 Color : COLOR0; }; struct VSOut { float4 Position : SV_Position; float2 Uv : TEXCOORD0; float4 Color : COLOR0; }; [shader("vertex")] VSOut VSMain(VSIn i) { VSOut o; o.Position = mul(ViewProjection, float4(i.Position, 0, 1)); o.Uv = i.Uv; o.Color = i.Color * Tint; return o; } [shader("pixel")] float4 PSMain(VSOut i) : SV_Target { return Atlas.Sample(AtlasSampler, i.Uv) * i.Color; } ``` The build generates a static class named after the file, `Sprite`, in the namespace `.Shaders`. `Program.cs`: ```csharp using System.Numerics; using System.Runtime.InteropServices; using Start.Shaders; // .Shaders // The SPIR-V of each entry point is embedded in the assembly. ReadOnlySpan vs = Sprite.VSMain.SpirV; ReadOnlySpan ps = Sprite.PSMain.SpirV; Console.WriteLine($"VSMain: {vs.Length} bytes of SPIR-V, PSMain: {ps.Length} bytes"); // The constant buffer as a C# struct, laid out like the shader's. var frame = new Sprite.Frame { ViewProjection = Matrix4x4.Identity, Tint = Vector4.One }; ReadOnlySpan upload = MemoryMarshal.AsBytes(new ReadOnlySpan(in frame)); Console.WriteLine($"Frame: {Sprite.Frame.Size} bytes, Tint at offset {Marshal.OffsetOf("Tint")}, {upload.Length} bytes to upload"); // Vertex input: one interleaved buffer. Console.WriteLine($"Vertex stride: {Sprite.VSMain.Stride}"); foreach (var a in Sprite.VSMain.Inputs) Console.WriteLine($" location {a.Location}: {a.Format} at offset {a.Offset}"); // Where each resource is bound. foreach (var b in Sprite.Bindings.All) Console.WriteLine($"{b.Name}: set {b.Set}, binding {b.Binding}, {b.Kind}, stages {b.Stages}"); ``` `dotnet run` prints: ```text VSMain: 1272 bytes of SPIR-V, PSMain: 856 bytes Frame: 80 bytes, Tint at offset 64, 80 bytes to upload Vertex stride: 20 location 0: Float32x2 at offset 0 location 1: Float32x2 at offset 8 location 2: Unorm8x4 at offset 16 Frame: set 0, binding 0, ConstantBuffer, stages Vertex Atlas: set 0, binding 1, SampledTexture, stages Pixel AtlasSampler: set 0, binding 2, Sampler, stages Pixel ``` - Entry points come from the `[shader("vertex")]`, `[shader("pixel")]` and `[shader("compute")]` attributes. Each becomes a nested class with `Stage`, `EntryPoint` and `SpirV`, a span over bytes embedded in the assembly. Nothing is read from disk at run time. - `Sprite.Frame` is a blittable struct with an explicit layout whose field offsets equal the compiled shader’s, so its bytes go to the GPU as they are. Vectors and matrices are `System.Numerics` types. A `Matrix4x4` member’s element (i, j) is HLSL’s `M[i][j]`. - `[[oneshader::format("unorm8x4")]]` declares a packed vertex format: the buffer holds 4 bytes, the shader sees a `float4`. That’s why the stride is 20, not 32. DXC compiles the same file with a warning that it ignores the attribute. - Without `[[vk::binding]]` or a register, resources are numbered in declaration order. Layouts and bindings follow the rules DXC uses for Vulkan, so existing shaders keep theirs. All the generated members are listed in [Generated bindings](https://recursiveai.org/oneshader/docs/generated-bindings/). ## Shader errors are build errors Misspell `Atlas` in the pixel shader and build again: ```text /tmp/oneshader-docs/Start/Shaders/Sprite.hlsl(24,45): error OSH0300: 'Atlass' is not declared. Did you mean 'Atlas'? [/tmp/oneshader-docs/Start/Start.csproj] ``` The error is on the HLSL line and column, with a stable `OSHnnnn` code, so the IDE jumps to it like a C# error. A shader is compiled again only when it, a file it includes, or the options change. ## Add Metal and WebGPU SPIR-V is the default. Select more targets in the project file: ```xml Exe net10.0 enable enable SpirV;Msl;Wgsl ``` With the same `Shaders/Sprite.hlsl`: ```csharp using System.Text; using Targets.Shaders; // WGSL for WebGPU and MSL for Metal, next to the SPIR-V. string wgsl = Encoding.UTF8.GetString(Sprite.PSMain.Wgsl); string msl = Encoding.UTF8.GetString(Sprite.PSMain.Msl); Console.WriteLine($"WGSL entry: {Sprite.PSMain.WgslEntryPoint}, MSL function: {Sprite.PSMain.MslFunctionName}"); Console.WriteLine($"Metal vertex buffer index: {Sprite.VSMain.MslVertexBufferIndex}"); Console.WriteLine($"Atlas: @group({Sprite.Bindings.Atlas.Set}) @binding({Sprite.Bindings.Atlas.Binding}), Metal texture({Sprite.Bindings.Atlas.MslIndex})"); Console.WriteLine(); Console.WriteLine(wgsl); File.WriteAllText("sprite.ps.metal", msl); ``` The start of the output: ```text WGSL entry: PSMain, MSL function: PSMain Metal vertex buffer index: 30 Atlas: @group(0) @binding(1), Metal texture(0) diagnostic(off, derivative_uniformity); struct PSMain_Output { @location(0) PSMain_result: vec4, } @group(0) @binding(1) var Atlas: texture_2d; @group(0) @binding(2) var AtlasSampler: sampler; ... @fragment fn PSMain(@builtin(position) in_PSMain_i_Position: vec4, @location(0) in_PSMain_i_Uv: vec2, @location(1) in_PSMain_i_Color: vec4) -> PSMain_Output { ... ``` - WGSL uses the Vulkan numbers as `@group` and `@binding`. Metal counts buffers, textures and samplers separately, so `Atlas` is texture 0. The interleaved vertex buffer is Metal buffer 30. - The constant-buffer and vertex structs are the same for every target: all three backends use one byte layout. - A shader that a target can’t express fails the build with an error that names the target, never a silent miscompile. Clip distances on WGSL, for example, are error OSH5202. That’s why the default is SPIR-V only. A target you don’t select costs no build time. - Clip-space Y points up in Metal and WebGPU and down in Vulkan. The shader isn’t adjusted, so flip it with the viewport. ## Build pipeline layouts `OneShader.Layouts` turns a shader into the layouts a renderer creates, as plain data with no graphics API behind it. In the project with all three targets: ```csharp using OneShader.Layouts; using OneShader.Layouts.Metal; using OneShader.Layouts.Vulkan; using OneShader.Layouts.WebGpu; using Layouts.Shaders; // The vertex and pixel entry points of Sprite.hlsl, merged. PipelineLayout pipeline = Sprite.Pipeline; // Vulkan: descriptor set layouts and vertex input, with the numeric Vk* enum values. PipelineLayoutInfo vulkan = pipeline.ToVulkan(); foreach (DescriptorSetLayoutInfo set in vulkan.SetLayouts) foreach (DescriptorSetLayoutBindingInfo b in set.Bindings) Console.WriteLine($"set {set.Set} binding {b.Binding}: {b.DescriptorType} ({(uint)b.DescriptorType}) x{b.DescriptorCount}, stages {b.StageFlags}"); VertexInputInfo input = Sprite.VSMain.Layout.Vertex!.ToVulkan(); foreach (VertexInputAttributeInfo a in input.Attributes) Console.WriteLine($"location {a.Location}: {a.Format} ({(uint)a.Format}) at {a.Offset}, stride {input.Bindings[0].Stride}"); // Metal: the argument index of each resource, per stage. if (pipeline.TryToMetal(out MetalPipelineLayout? metal, out var metalError)) foreach (MetalStageLayout stage in metal!.Stages) foreach (MetalArgument arg in stage.Arguments) Console.WriteLine($"Metal {stage.Stage}: {arg.Name} -> {arg.Kind}({arg.Index})"); else Console.WriteLine(metalError); // WebGPU: bind group layout entries, with the API's strings. if (pipeline.TryToWebGpu(out WebGpuPipelineLayout? webgpu, out var webgpuError)) foreach (BindGroupLayoutInfo group in webgpu!.Groups) foreach (BindGroupLayoutEntryInfo e in group.Entries) Console.WriteLine($"WebGPU group {group.Group} binding {e.Binding}: {e.Name}, {e.Kind}, visibility {e.Visibility}"); else Console.WriteLine(webgpuError); ``` ```text set 0 binding 0: UniformBuffer (6) x1, stages Vertex set 0 binding 1: SampledImage (2) x1, stages Fragment set 0 binding 2: Sampler (0) x1, stages Fragment location 0: R32G32Sfloat (103) at 0, stride 20 location 1: R32G32Sfloat (103) at 8, stride 20 location 2: R8G8B8A8Unorm (37) at 16, stride 20 Metal Vertex: Frame -> Buffer(0) Metal Pixel: Atlas -> Texture(0) Metal Pixel: AtlasSampler -> Sampler(0) WebGPU group 0 binding 0: Frame, Buffer, visibility Vertex WebGPU group 0 binding 1: Atlas, Texture, visibility Fragment WebGPU group 0 binding 2: AtlasSampler, Sampler, visibility Fragment ``` - The numbers in parentheses are the values of `VkDescriptorType` and `VkFormat`, ready to cast into whichever Vulkan binding you use. - `Sprite.Pipeline` is the file’s vertex and pixel entry points (or its one compute entry point). For other combinations, `PipelineLayout.TryCreate` merges any entry layouts and returns a conflict as a `LayoutError` value, never an exception. - `TryToMetal` fails with `MissingTarget` when `Msl` wasn’t built. `TryToWebGpu` reports what WebGPU can’t express, such as resource arrays, with what to change. - Creating layouts allocates, so do it when a shader loads, not every frame. ## Compile at run time For tools, editors and hot reload, call the compiler directly. It’s the `RecursiveAI.OneShader` package (`dotnet add package RecursiveAI.OneShader --version 1.2.0`), namespace `OneShader`: ```csharp using OneShader; using OneShader.Reflection; CompileResult result = ShaderCompiler.CompileFile("Shaders/Sprite.hlsl", new CompileOptions { Targets = [CompileTarget.SpirV, CompileTarget.Wgsl], }); foreach (var d in result.Diagnostics) Console.WriteLine(d); // Shaders/Sprite.hlsl(24,45): error OSH0300: ... if (!result.Success) return 1; foreach (ShaderBlob blob in result.Blobs) Console.WriteLine($"{blob.EntryPoint} ({blob.Stage}) -> {blob.Target}: {blob.Data.Length} bytes"); File.WriteAllBytes("sprite.ps.spv", result.Find("PSMain", CompileTarget.SpirV)!.ToArray()); File.WriteAllText("sprite.ps.wgsl", result.Find("PSMain", CompileTarget.Wgsl)!.AsText()); // Reflection: what the shader binds, where, and in which stages. foreach (ResourceInfo r in result.Reflection!.Resources) Console.WriteLine($"{r.Name}: {r.Kind}, set {r.Slots[0].Set} binding {r.Slots[0].Binding}, {r.Stages}"); return 0; ``` ```text VSMain (Vertex) -> SpirV: 1272 bytes VSMain (Vertex) -> Wgsl: 1640 bytes PSMain (Pixel) -> SpirV: 856 bytes PSMain (Pixel) -> Wgsl: 841 bytes Frame: ConstantBuffer, set 0 binding 0, Vertex Atlas: SampledTexture, set 0 binding 1, Pixel AtlasSampler: Sampler, set 0 binding 2, Pixel ``` - Errors come back as diagnostics with file, line and column, never as exceptions: bad input gives diagnostics, not an exception or a stack overflow. `ShaderCompiler.Compile(source, path, options)` takes the text instead of a path. - Output is deterministic: the same source and options give the same bytes. `Compile` can be called from several threads at once. - `sprite.ps.spv` passes `spirv-val` for Vulkan 1.2, and `sprite.ps.wgsl` passes naga. - `result.Files` lists every file the compile read, including `#include`s, for a file watcher. See [Reload shaders while the app runs](https://recursiveai.org/oneshader/cookbook/hot-reload-shaders/). ## Compile from the command line `RecursiveAI.OneShader.Tool` is the `oneshader` command, with DXC’s core flags: ```sh dotnet tool install -g RecursiveAI.OneShader.Tool --version 1.2.0 oneshader -T ps_6_0 -E PSMain -spirv Sprite.hlsl -Fo sprite.ps.spv oneshader --all-entries --target spirv --target msl --target wgsl --reflect sprite.json Sprite.hlsl -Fo out/sprite ``` The second command writes `sprite.VSMain.spv`, `.metal` and `.wgsl` and the same three for `PSMain`, plus the reflection as JSON. The tool needs the .NET 10 runtime. Exit code 0 is success, 1 a shader that didn’t compile, 2 a wrong command line. See [Switch from the DXC command line](https://recursiveai.org/oneshader/cookbook/migrate-from-dxc/). ## New in 1.2.0 - **Shader variants in the build.** `Axes` metadata on an item (such as `ALPHA_TEST` and `LIGHTING=NONE|LAMBERT|PBR`) compiles every combination, and the generated `Key` picks one at run time. See [Shader variants with Axes](https://recursiveai.org/oneshader/cookbook/shader-variants-axes/). - **Specialization constants.** `[[vk::constant_id(N)]] const` values set when the pipeline is created, on Vulkan, Metal and WebGPU, with a generated `Specialization` struct. See [Specialization constants](https://recursiveai.org/oneshader/cookbook/specialization-constants/). - **16-bit and 64-bit types.** `half`, `int16_t`, `double`, `uint64_t` and 64-bit atomics, as DXC compiles them, with the device features each entry point needs. See [16-bit and 64-bit types](https://recursiveai.org/oneshader/cookbook/hlsl-16-bit-64-bit-types/). - **A compile cache** for run-time compiles. See [Cache shader compiles](https://recursiveai.org/oneshader/cookbook/shader-compile-cache/). The full list is in the [changelog](https://recursiveai.org/oneshader/docs/changelog/). ## The packages All three packages share one version number. | Package | What it is | License | |---|---|---| | `RecursiveAI.OneShader.Build` | MSBuild task and source generator: compiles `*.hlsl` during `dotnet build`, reports errors on the HLSL line, embeds the selected targets, builds shader variants and generates C# structs, bindings and layouts. Guide: [the build package](https://recursiveai.org/oneshader/docs/build-package/). Brings `RecursiveAI.OneShader` with it. | MIT or Apache-2.0 | | `RecursiveAI.OneShader` | The compiler: `ShaderCompiler`, reflection, `InterfaceCompatibility` for hot reload, `ShaderCompileCache`, pipeline layouts and specialization values (`OneShader.Layouts`), and the support types the generated code uses. Namespace `OneShader`. Guide: [the compiler package](https://recursiveai.org/oneshader/docs/compiler/). No `unsafe` code, NativeAOT and trimming supported. | MIT or Apache-2.0 | | `RecursiveAI.OneShader.Tool` | The `oneshader` command-line compiler, a .NET tool with DXC-style flags, for CMake, Make and scripts. Guide: [the command line](https://recursiveai.org/oneshader/docs/command-line/). | MIT or Apache-2.0 | ## What OneShader 1.2.0 doesn’t do - **Other shader stages.** Vertex, pixel and compute only: no geometry, tessellation, mesh or ray-tracing shaders. - **Direct3D bytecode.** There’s no DXIL output. For Direct3D 12, keep DXC. - **GLSL.** The input is HLSL and the outputs are SPIR-V, MSL and WGSL: GLSL is neither read nor written. - **Unbounded resource arrays** (`Texture2D Tex[]`) and the partitioned wave operations `WaveMatch` and `WaveMultiPrefix*` are errors (OSH0350). Bounded arrays and the other wave and quad intrinsics work. - **Every feature on every target.** Some constructs exist on one target and not another: `double`, for example, is SPIR-V only. The [MSL](https://recursiveai.org/oneshader/docs/msl-backend/) and [WGSL](https://recursiveai.org/oneshader/docs/wgsl-backend/) pages list them, and each one is a compile error naming the target. ## Where next Cookbook recipes, each with a complete, run sample: - [Compile HLSL to SPIR-V for Vulkan](https://recursiveai.org/oneshader/cookbook/hlsl-to-spirv/), [convert HLSL to WGSL](https://recursiveai.org/oneshader/cookbook/hlsl-to-wgsl/), [compile HLSL to Metal Shading Language](https://recursiveai.org/oneshader/cookbook/hlsl-to-metal/) - [Compile HLSL in a C# project with dotnet build](https://recursiveai.org/oneshader/cookbook/compile-hlsl-csharp/), [match constant buffers with C# structs](https://recursiveai.org/oneshader/cookbook/cbuffer-packing-csharp-structs/), [push constants on every target](https://recursiveai.org/oneshader/cookbook/push-constants/) - [Read shader reflection in C#](https://recursiveai.org/oneshader/cookbook/spirv-reflection-csharp/), [reload shaders while the app runs](https://recursiveai.org/oneshader/cookbook/hot-reload-shaders/), [shader variants with defines and includes](https://recursiveai.org/oneshader/cookbook/shader-variants-defines-includes/) - [Write a compute shader for Vulkan and WebGPU](https://recursiveai.org/oneshader/cookbook/hlsl-compute-shader/), [wave intrinsics on Vulkan, Metal and WebGPU](https://recursiveai.org/oneshader/cookbook/hlsl-wave-intrinsics/), [InterlockedAdd and float atomics](https://recursiveai.org/oneshader/cookbook/hlsl-interlockedadd/), [the precise keyword](https://recursiveai.org/oneshader/cookbook/hlsl-precise-keyword/) - [Shader variants with Axes](https://recursiveai.org/oneshader/cookbook/shader-variants-axes/), [specialization constants](https://recursiveai.org/oneshader/cookbook/specialization-constants/), [16-bit and 64-bit types](https://recursiveai.org/oneshader/cookbook/hlsl-16-bit-64-bit-types/), [cache shader compiles](https://recursiveai.org/oneshader/cookbook/shader-compile-cache/) - [Switch from the DXC command line](https://recursiveai.org/oneshader/cookbook/migrate-from-dxc/) The reference docs: - The [overview](https://recursiveai.org/oneshader/docs/overview/) and the package guides: [build](https://recursiveai.org/oneshader/docs/build-package/), [compiler](https://recursiveai.org/oneshader/docs/compiler/), [command line](https://recursiveai.org/oneshader/docs/command-line/) - [Generated bindings](https://recursiveai.org/oneshader/docs/generated-bindings/) and the [reflection file format](https://recursiveai.org/oneshader/docs/reflection-format/) - The backends: [SPIR-V](https://recursiveai.org/oneshader/docs/spirv-backend/), [Metal Shading Language](https://recursiveai.org/oneshader/docs/msl-backend/), [WGSL](https://recursiveai.org/oneshader/docs/wgsl-backend/), and the [intermediate representation](https://recursiveai.org/oneshader/docs/ir/) they share - How it’s checked and how fast it is: [WebGPU execution parity](https://recursiveai.org/oneshader/docs/webgpu-parity/), [performance against DXC](https://recursiveai.org/oneshader/docs/performance/), and the [changelog](https://recursiveai.org/oneshader/docs/changelog/) Guides: [HLSL vs GLSL](https://recursiveai.org/oneshader/articles/hlsl-vs-glsl/) and [HLSL to SPIR-V compiler compared](https://recursiveai.org/oneshader/articles/hlsl-to-spirv-compiler/). # OneShader **A HLSL compiler for .NET, written in C# with no dependency beyond .NET itself.** Modern HLSL (Shader Model 6 style) for vertex, pixel and compute shaders goes in; SPIR-V for Vulkan, Metal Shading Language and WGSL come out, each with reflection. An MSBuild task and a source generator compile your `*.hlsl` files during `dotnet build`, report errors on the HLSL line, and generate C# structs and binding code that match the compiled shader. [![NuGet](https://img.shields.io/nuget/v/RecursiveAI.OneShader)](https://www.nuget.org/packages/RecursiveAI.OneShader) [![CI](https://github.com/recursive-ai-oss/oneshader-csharp/actions/workflows/ci.yml/badge.svg)](https://github.com/recursive-ai-oss/oneshader-csharp/actions/workflows/ci.yml) [![.NET 10](https://img.shields.io/badge/.NET-10.0-512BD4)](#install) [![License](https://img.shields.io/badge/license-MIT%20OR%20Apache--2.0-blue)](#license) How many times faster or slower OneShader compiles a shader than DXC (the released binary, SPIR-V output), one process per entry point over 68 corpus entry points, geometric mean, log scale. A bar to the right is a win for OneShader, a bar to the left is a loss: Bar chart: how many times faster or slower OneShader compiles than DXC, log scale. The NativeAOT build is about 2.4 times faster when warm and 2.7 times faster on the first run. Started as an ordinary .NET process, it is slower: about 5 to 6 times with ReadyToRun and 13 to 14 times with the JIT. [Full results, with allocations per compile, output size, the losses and the method →](https://recursiveai.org/oneshader/docs/performance/) ## Why OneShader - **Shader errors are build errors.** Put `*.hlsl` in a project and reference `RecursiveAI.OneShader.Build`: a mistake shows up as `Shaders/Sprite.hlsl(24,45): error OSH0300: 'Atlass' is not declared. Did you mean 'Atlas'?` on the HLSL line, with a stable code, not as a black screen at run time. - **C# that matches the shader.** Constant buffers, vertex inputs and resource bindings become blittable structs and constants whose offsets equal the compiled shader's, so a CPU/GPU layout mismatch fails the build ([generated bindings](https://recursiveai.org/oneshader/docs/generated-bindings/)). One layout serves every target. - **Three targets from one front end.** SPIR-V (Vulkan 1.2+) is the default; Metal Shading Language and WGSL are opt-in per project (`SpirV;Msl;Wgsl`). A feature a target cannot express is a diagnostic naming the target, never a silent miscompile ([MSL](https://recursiveai.org/oneshader/docs/msl-backend/), [WGSL](https://recursiveai.org/oneshader/docs/wgsl-backend/)). - **Checked against DXC and real GPUs.** SPIR-V is validated with `spirv-val` and every corpus shader is run on the GPU next to the same shader compiled by DXC, with the Vulkan validation layers on; WGSL goes through `naga` and a browser WebGPU run ([WebGPU parity](https://recursiveai.org/oneshader/docs/webgpu-parity/)); MSL goes through Apple's compiler and a Metal run. - **Safe on text you did not write.** The compiler runs inside your build, so bad input yields diagnostics, never an exception or a stack overflow, and the front end is fuzzed ([fuzzing](https://github.com/recursive-ai-oss/oneshader-csharp/blob/8896e63be3a4f246295427451f15e4fcfb5535cf/fuzz/README.md)). It contains no `unsafe` code, uses no reflection, and is NativeAOT and trim clean. - **Fast, with the losses shown.** As a NativeAOT binary it compiles the corpus about 2.4 times faster than DXC, warm, and every one of the 68 entry points is faster. Run as an ordinary .NET process it is slower than DXC on every one of them (about 6 times with ReadyToRun, 14 times with the JIT), because starting the runtime and the compiler's code dominates a compile that takes 0.34 ms at the median once the process is running. Its SPIR-V is 5.7% larger than DXC's by default because it keeps debug names, and 8.8% smaller with them off. Memory use was not measured ([every result](https://recursiveai.org/oneshader/docs/performance/)). ## Install For building shaders into your project: ```sh dotnet add package RecursiveAI.OneShader.Build ``` For the compiler itself, for hot reload and tools: ```sh dotnet add package RecursiveAI.OneShader ``` From a command line or another build system (CMake, Make, scripts), with a DXC-style interface: ```sh dotnet tool install -g RecursiveAI.OneShader.Tool # then: oneshader -T ps_6_0 -E PSMain -spirv shader.hlsl -Fo shader.spv ``` Target framework `net10.0`. Build with the .NET SDK (`dotnet build`); Visual Studio's .NET Framework MSBuild cannot load the build task. Package guides: [RecursiveAI.OneShader.Build](https://recursiveai.org/oneshader/docs/build-package/), [RecursiveAI.OneShader](https://recursiveai.org/oneshader/docs/compiler/) and [RecursiveAI.OneShader.Tool](https://recursiveai.org/oneshader/docs/command-line/). ## Quick start `Shaders/Sprite.hlsl`: ```hlsl cbuffer Frame : register(b0) { float4x4 ViewProjection; }; [[vk::binding(1)]] Texture2D Atlas; [[vk::binding(2)]] SamplerState AtlasSampler; struct VSIn { float2 Position : POSITION; float2 Uv : TEXCOORD0; float4 Color : COLOR0; }; struct VSOut { float4 Position : SV_Position; float2 Uv : TEXCOORD0; float4 Color : COLOR0; }; [shader("vertex")] VSOut VSMain(VSIn i) { VSOut o; o.Position = mul(ViewProjection, float4(i.Position, 0, 1)); o.Uv = i.Uv; o.Color = i.Color; return o; } [shader("pixel")] float4 PSMain(VSOut i) : SV_Target { return Atlas.Sample(AtlasSampler, i.Uv) * i.Color; } ``` After `dotnet build`, with no reflection and no file I/O at run time: ```csharp using MyGame.Shaders; // .Shaders ReadOnlySpan vs = Sprite.VSMain.SpirV; // embedded SPIR-V var frame = new Sprite.Frame { ViewProjection = camera }; // explicit layout, Size = 64 uint set = Sprite.Bindings.Atlas.Set, binding = Sprite.Bindings.Atlas.Binding; ``` To compile at run time instead (tools, hot reload): ```csharp using OneShader; CompileResult result = ShaderCompiler.CompileFile("Shaders/Sprite.hlsl"); if (!result.Success) { foreach (var d in result.Diagnostics) Console.WriteLine(d); return; } byte[] spirv = result.Find("PSMain", CompileTarget.SpirV)!.ToArray(); ``` ## Status **Version 1.2.** Vertex, pixel and compute shaders compile to optimized SPIR-V, Metal Shading Language and WGSL with reflection; the build package embeds each selected target next to generated C# bindings (the `OneShaderTargets` switch), and the runtime API supports hot reload with a layout-compatibility check. Not supported: geometry, tessellation, mesh and ray-tracing stages, unbounded resource arrays and the SM 6.5 partitioned wave operations (`WaveMatch`, `WaveMultiPrefix*`), and some features on some targets (the backend pages list them). The public API and the generated code follow [Semantic Versioning](https://semver.org/); the [changelog](https://recursiveai.org/oneshader/docs/changelog/) lists what each release changes. ## Where to go - **Using it:** the package guides ([build](https://recursiveai.org/oneshader/docs/build-package/), [compiler](https://recursiveai.org/oneshader/docs/compiler/)), the [generated bindings](https://recursiveai.org/oneshader/docs/generated-bindings/), and the [reflection format](https://recursiveai.org/oneshader/docs/reflection-format/). - **Backends:** [SPIR-V](https://recursiveai.org/oneshader/docs/spirv-backend/), [MSL](https://recursiveai.org/oneshader/docs/msl-backend/), [WGSL](https://recursiveai.org/oneshader/docs/wgsl-backend/), the shared [IR](https://recursiveai.org/oneshader/docs/ir/), and [browser WebGPU parity](https://recursiveai.org/oneshader/docs/webgpu-parity/). - **Evidence:** [performance and its losses](https://recursiveai.org/oneshader/docs/performance/) (the numbers behind the chart; reproduce with `bench/run.sh`), [fuzzing](https://github.com/recursive-ai-oss/oneshader-csharp/blob/8896e63be3a4f246295427451f15e4fcfb5535cf/fuzz/README.md), and the corpus in [`corpus/`](https://github.com/recursive-ai-oss/oneshader-csharp/blob/8896e63be3a4f246295427451f15e4fcfb5535cf/corpus/README.md). - **What changed:** [CHANGELOG.md](https://recursiveai.org/oneshader/docs/changelog/). ## Building ```sh dotnet build dotnet test ``` Build output goes to the directory named by `ONESHADER_ARTIFACTS` (or `.artifacts` in the repository when it is unset). Tests that compare against reference compilers use released binaries fetched by `tools/fetch-oracles.sh`; versions are pinned in `tools/oracles.env`. ## License Licensed under either of [MIT](https://github.com/recursive-ai-oss/oneshader-csharp/tree/8896e63be3a4f246295427451f15e4fcfb5535cf/LICENSE-MIT) or [Apache-2.0](https://github.com/recursive-ai-oss/oneshader-csharp/tree/8896e63be3a4f246295427451f15e4fcfb5535cf/LICENSE-APACHE) at your option. Copyright (c) Recursive AI, LLC. # OneShader.Build **HLSL shaders compiled during `dotnet build`, with C# bindings generated from them.** Put `*.hlsl` files in a project and reference this package: each shader is compiled to SPIR-V by [OneShader](https://www.nuget.org/packages/RecursiveAI.OneShader), a HLSL compiler written in C# with no dependency beyond .NET itself. Shader errors are build errors on the HLSL line, and constant buffers, vertex inputs and resource bindings become C# types and constants laid out to match the compiled shader, so a CPU/GPU mismatch fails the build instead of the frame. Vulkan (SPIR-V) is the default target; Metal (MSL) and WebGPU (WGSL) are opt-in with `OneShaderTargets`. ## Install ```sh dotnet add package RecursiveAI.OneShader.Build ``` Target framework `net10.0`. Build with the .NET SDK (`dotnet build`); Visual Studio's .NET Framework MSBuild cannot load the build task. ## Quick start `Shaders/Sprite.hlsl`: ```hlsl cbuffer Frame : register(b0) { float4x4 ViewProjection; }; [[vk::binding(1)]] Texture2D Atlas; [[vk::binding(2)]] SamplerState AtlasSampler; struct VSIn { float2 Position : POSITION; float2 Uv : TEXCOORD0; [[oneshader::format("unorm8x4")]] float4 Color : COLOR0; }; struct VSOut { float4 Position : SV_Position; float2 Uv : TEXCOORD0; float4 Color : COLOR0; }; [shader("vertex")] VSOut VSMain(VSIn i) { VSOut o; o.Position = mul(ViewProjection, float4(i.Position, 0, 1)); o.Uv = i.Uv; o.Color = i.Color; return o; } [shader("pixel")] float4 PSMain(VSOut i) : SV_Target { return Atlas.Sample(AtlasSampler, i.Uv) * i.Color; } ``` After `dotnet build`, with no reflection and no file I/O at run time: ```csharp using MyGame.Shaders; // .Shaders ReadOnlySpan vs = Sprite.VSMain.SpirV; // embedded SPIR-V ReadOnlySpan ps = Sprite.PSMain.SpirV; uint stride = Sprite.VSMain.Stride; // 20 foreach (var a in Sprite.VSMain.Inputs) { /* a.Location, a.Format (Float32x2, Unorm8x4, ...), a.Offset */ } var frame = new Sprite.Frame { ViewProjection = camera }; // explicit layout, Size = 64 uint set = Sprite.Bindings.Atlas.Set, binding = Sprite.Bindings.Atlas.Binding; ``` A mistake in the shader is an ordinary build error: ``` Shaders/Sprite.hlsl(24,45): error OSH0300: 'Atlass' is not declared. Did you mean 'Atlas'? ``` ## What is generated One static class per HLSL file (named after the file, in PascalCase), containing: - one blittable struct per constant buffer, with field offsets equal to the compiled shader's; matrices are `Matrix4x4` members where element (i, j) is HLSL's `M[i][j]`; - a struct per structured-buffer element type; - `Bindings`: set, binding, count and kind of every resource, as constants; - for a `[[vk::push_constant]]` block: a blittable `PushConstants` struct and `Bindings.PushConstants` with its `Size`, the stages for the `VkPushConstantRange`, and the Metal index and WebGPU group and binding of the targets you selected; - per entry point: `Layout`, its resources, push-constant block and vertex input as plain data, and per file `Pipeline` (vertex + pixel or one compute entry point) with `ToVulkan()`, `TryToMetal()` and `TryToWebGpu()` giving descriptor set layouts, push-constant ranges, Metal argument indices and WebGPU bind group layouts (see the `RecursiveAI.OneShader` readme, "Pipeline layouts"); - for `[[vk::constant_id]]` specialization constants: `SpecializationConstants..Id`/`Default`, a blittable `Specialization` struct with `Defaults` and `ToVulkan()`, `ToMetal()` and `ToWebGpu()`, and `SpecConstantIds` on the entry points that read them (see docs/generated-bindings.md); - per entry point: stage, entry name, the SPIR-V bytes; for vertex shaders a `Vertex` struct, `Stride` and `Inputs`; for compute shaders the thread-group size. Entry points also carry the MSL and WGSL sources when those targets are selected (next section). ## Choosing targets ```xml SpirV;Msl;Wgsl ``` `OneShaderTargets` is a `;`-list of `SpirV`, `Msl` and `Wgsl` (case-insensitive). A `OneShader` item can override it with `Targets` metadata. The default stays SPIR-V only because some valid shaders are refused by a target (for example clip distances or resource arrays on WGSL, OSH5202); a refusal is a build error on the HLSL line, `Shader.hlsl(12,5): error OSH5202: ...`, only when that target is selected, and a target nobody asked for costs no build time. With a target selected, every entry point class gains (in addition to `SpirV`): | Member | Target | Meaning | |---|---|---| | `Msl` | Msl | `ReadOnlySpan`, UTF-8 Metal Shading Language source | | `MslFunctionName` | Msl | the function to pass to `newFunction(name:)`; differs from `EntryPoint` when MSL reserves the name (`main` becomes `main0`) | | `MslVertexBufferIndex` | Msl | vertex shaders with inputs: the Metal buffer index the interleaved vertex buffer (`Stride`, `Inputs`) is bound at | | `Wgsl` | Wgsl | `ReadOnlySpan`, UTF-8 WGSL source | | `WgslEntryPoint` | Wgsl | the entry function name in the WGSL source | Bindings carry the numbers of every selected target. `Set` and `Binding` are Vulkan's (WGSL uses the same numbers as `@group`/`@binding`); with Msl selected, each resource class also has `MslIndex` (and `MslCounterIndex` for Append/Consume counters), the Metal argument index counted within the resource's class (buffer, texture, sampler) in declaration order, and `ResourceBinding.MetalIndex` in `Bindings.All` holds the same number (`ResourceBinding.NoMetalIndex` when Msl is not selected). With only Msl selected, `Set` is 0 and `Binding` repeats the Metal index. Constant-buffer and vertex structs are the same for every target: all backends use one byte layout. Support types (`ShaderStage`, `ResourceKind`, `VertexFormat`, `VertexAttribute`, `ResourceBinding`) come from `RecursiveAI.OneShader`, so a rendering layer in another assembly can take them. When only these are used, trimming removes the compiler from a NativeAOT app. Item metadata: `EntryPoints`, `Defines`, `IncludeDirectories`, `Namespace`, `ClassName`, `Targets`, `Axes`, `MaxVariants`. Properties: `EnableDefaultOneShaderItems`, `OneShaderNamespace`, `OneShaderTargets`, `OneShaderEnable16BitTypes` (item metadata `Enable16BitTypes`), `OneShaderDefines`, `OneShaderIncludeDirectories`, `OneShaderMaxVariants`. Shaders are recompiled only when they, a file they include, or the options change; the output is deterministic. ## Variants One HLSL file can be compiled many times with different macros, for a material system that picks the right shader at run time. Two shapes, both in the project file: ```xml ``` Listing a file as an item gives it metadata; the default `**/*.hlsl` glob skips files that are already listed. - **Axes** is a `;`-list. A bare `NAME` is a flag: the variant either defines `NAME` as `1` or does not define it (`#if ALPHA_TEST`). `NAME=A|B|C` is a choice: the variant defines `NAME` as the value *and* `NAME_VALUE` as `1` (`LIGHTING=PBR` and `LIGHTING_PBR=1`, so `#if LIGHTING_PBR` works for words; numeric values also work with `#if LIGHTING == 2`). Names and values are letters, digits and `_`. The cartesian product is compiled, first axis varying slowest. `Defines`, `$(OneShaderDefines)` and the other item metadata apply to every variant. - **OneShaderVariant** items name a shader (`Shader`, any path to the same file) and give a `Defines` list. Use either `Axes` or `OneShaderVariant` items on one shader, not both. - **Limit.** A shader may have 64 variants; more is error `OSH0905` naming the item and the fix. Raise it with `MaxVariants="256"` metadata or `` in the project. Each variant is compiled for every selected target and embedded in the assembly, so the limit is also a size guard. - **One compile per variant per target**, in parallel. A variant is recompiled only when its define set, the shader, a file it includes or the options change: adding a value to an axis compiles just the new combinations. Each variant's SPIR-V is byte-for-byte what `ShaderCompiler.Compile` produces with the same defines. - **Errors** are the usual `Shaders/Material.hlsl(31,9): error OSH0300: ...` on the HLSL line, followed by the define set it happened under, `(variant ALPHA_TEST=1, LIGHTING=PBR)`. The same error under several variants is reported once, with `(variants A; B; C and 2 more)` or `(in all 6 variants)`. Generated code for a file with variants (`Material.hlsl` above, namespace and class as for any file): ```csharp var v = Material.Variants.Get(new Material.Key(alphaTest: true, lighting: Material.LightingValue.Pbr)); ReadOnlySpan spirv = v.PSMain.SpirV; // Msl / Wgsl / MslFunctionName / Layout / ThreadGroupSizeX ... as for a single file ReadOnlySpan bindings = v.Bindings; // can differ per variant (the normal map exists only when lit) PipelineLayout? pipeline = v.Pipeline; ``` - `Material.Key` is a struct with one property per axis (`bool` for flags, a generated enum such as `LightingValue` for choices; `VariantName` for `OneShaderVariant` sets). `Variants.Get(key)` is an array index: no allocation, no reflection; it throws for an enum value that is out of range, `TryGet` does not. `Variants.All` and `Variants.Count` list every variant; `Variant.Name`, `Variant.Defines` and `Variant.Index` identify one. - The `Variant` has one property per entry point (`PSMain`, null in variants that lack it; non-null when every variant has it). Each is an `Entry` with `SpirV`, `Msl`, `Wgsl` (empty when the target was not selected), `MslFunctionName`, `WgslEntryPoint`, `Stride`/`Inputs` for vertex entries, `ThreadGroupSize*` for compute, `Layout`. - The same data is also available the way a single-variant file has it, one static class per variant: `Material.Variants.AlphaTestOn_LightingPbr.PSMain.SpirV`, `.Bindings.All`, `.Reflection`, `.Pipeline`. - **Shared structs.** A constant-buffer, push-constant or vertex struct that every variant agrees on is generated once at file level (`Material.Frame`, `Material.VSMainVertex`), so one CPU struct serves all variants. When a struct's layout differs between variants (a `cbuffer` member under `#if`), no shared type is generated: each variant's class gets its own (`Variants..Frame`) so a variant never gets a type that does not match its shader. Resource bindings can differ per variant and are read from the variant (`Variant.Bindings`, `Variant.Pipeline`). - A file without `Axes` or variant items generates exactly the same C# as before. ## What works - Vertex, pixel and compute shaders in modern HLSL (Shader Model 6 style, HLSL 2021): structs, functions, templates, member functions, inheritance and operator overloads, constant buffers, textures and samplers (including comparison sampling, gathers and LOD queries), structured, byte-address and Append/Consume buffers, storage images, bounded resource arrays (constant and run-time indexing), groupshared memory, barriers and atomics, and the preprocessor. - SPIR-V for Vulkan 1.2, optimized (inlining, SSA, constant folding, dead-code removal), validated with `spirv-val`, and checked by running every shader in its test corpus on a GPU against the same shader compiled by DXC: 2D sprites, text and UI, a PBR forward renderer with clustered lights, cascaded shadows, SSAO, bloom, tone mapping, image-based lighting and skinning, and compute culling. - Buffer layouts and descriptor bindings follow the conventions DXC uses for Vulkan, so existing shaders keep their layouts. - Metal Shading Language and WGSL embedded next to the SPIR-V when selected with `OneShaderTargets`, with their entry names and Metal binding indices in the generated code; the same corpus runs on Metal and in a browser on WebGPU against the DXC reference. - Push constants (`[[vk::push_constant]]`) on every target: a Vulkan push-constant block, a Metal `constant T&` buffer argument, a WebGPU uniform buffer; one set of bytes for all three. - Hot reload: generated classes carry the reflection they were built from, and `InterfaceCompatibility.Compare` tells whether a shader recompiled at run time still matches them (see the `RecursiveAI.OneShader` package). ## Known limits - Clip-space Y points up in Metal and WebGPU and down in Vulkan; the shader is not adjusted, so flip with the viewport. - Unbounded resource arrays and `WaveMatch`/`WaveMultiPrefix*` are not supported. Wave and quad intrinsics work; a shader that uses them needs subgroup support on the device (`SubgroupOperations` on the generated entry class and `Layout`). - `dotnet build` only (not Visual Studio's .NET Framework MSBuild). ## License MIT OR Apache-2.0, at your option. Copyright © Recursive AI, LLC. # OneShader **A HLSL compiler for .NET, written in C#, with no dependency beyond .NET itself.** Modern HLSL (Shader Model 6 style) for vertex, pixel and compute shaders goes in; SPIR-V for Vulkan comes out, with reflection: inputs and outputs, resource bindings, buffer layouts and thread-group sizes. NativeAOT and trimming are supported. Compiles to SPIR-V (Vulkan), Metal Shading Language and WGSL (WebGPU). Most projects want [RecursiveAI.OneShader.Build](https://www.nuget.org/packages/RecursiveAI.OneShader.Build), which compiles `*.hlsl` during `dotnet build` and generates C# bindings. This package is the compiler itself, for compiling at run time (hot reload, tools) and for the support types the generated code uses. ## Install ```sh dotnet add package RecursiveAI.OneShader ``` Target framework `net10.0`. The namespace is `OneShader`. ## Quick start ```csharp using OneShader; string source = File.ReadAllText("Shaders/Blur.hlsl"); CompileResult result = ShaderCompiler.Compile(source, "Shaders/Blur.hlsl", new CompileOptions { Defines = [new("RADIUS", "4")], IncludeDirectories = ["Shaders/Common"], }); foreach (var d in result.Diagnostics) Console.WriteLine(d); // Shaders/Blur.hlsl(12,5): error OSH0300: ... if (result.Success) { ReadOnlySpan spirv = result.Find("PSMain", CompileTarget.SpirV)!.Data; var reflection = result.Reflection!; // resources, bindings, buffer layouts, entry points } ``` Errors are returned as diagnostics with file, line and column, never thrown. Includes can come from disk or from memory through `CompileOptions.IncludeResolver`. Output is deterministic. ## Hot reload During development, watch the HLSL file, recompile it with the runtime API and swap the new SPIR-V in, as long as the interface the C# was built against has not changed. Projects using the build package get that interface as `Shaders.Sprite.Reflection` (generated from `Sprite.hlsl`). ```csharp CompileResult result = ShaderCompiler.CompileFile("Shaders/Sprite.hlsl"); // diagnostics, not exceptions watcher.Watch(result.Files); // the file and its #includes, even after an error if (!result.Success) { Print(result.Diagnostics); return; } // keep running the last good shader CompatibilityReport report = InterfaceCompatibility.Compare(Shaders.Sprite.Reflection, result.Reflection!); if (report.IsCompatible) renderer.Replace(result.Find("PSMain", CompileTarget.SpirV)!.Data); else Console.WriteLine(report); // "Frame.Tint moved from offset 64 to 80" ... "Rebuild the project ..." ``` Compatible means everything the generated C# relies on is unchanged: constant-buffer and structured-buffer layouts (member offsets, types, array and matrix strides, sizes), vertex input locations, formats, offsets and stride, resource set, binding, kind, array size and stage visibility, the entry points present and their stages, and compute thread-group sizes. Edits to function bodies, new entry points, and resources that become visible to fewer stages are compatible. `ShaderCompiler.Compile` can be called from several threads at once. `samples/HotReload` is a complete console loop. ## Compile cache `ShaderCompileCache` answers a compile from memory when the same inputs were compiled before, for hosts that build many variants of the same shaders (hot reload, tools, material systems). A result is reused when the source text, path, every option (targets, defines, include directories, entry points, matrix packing, 16-bit types, optimization, names) and the content of every include the earlier compile read are unchanged; includes are looked up again on each call, so an edited file, or a new file earlier on the search path, makes the next call compile. ```csharp var cache = new ShaderCompileCache(maxEntries: 256); // thread-safe; least recently used results are dropped CompileResult a = cache.Compile(source, "Material.hlsl", new CompileOptions { Defines = [new("ALPHA_TEST", null)] }); CompileResult b = cache.CompileFile("Shaders/Material.hlsl", options); // reads the file every call Console.WriteLine($"{cache.Hits} hits, {cache.Misses} compiles"); cache.Clear(); ``` Callers asking for one key at the same time share one compile; different keys compile in parallel. Results (errors included) are shared objects, so treat them as read-only; they are the bytes a direct `ShaderCompiler.Compile` gives. A cancelled call caches nothing. ## Pipeline layouts `OneShader.Layouts` turns a shader into what a renderer creates, as plain data with no graphics API behind it: descriptor set layouts and push-constant ranges for Vulkan, per-stage argument indices for Metal, bind group layouts for WebGPU, and vertex input for all three. It reads the generated classes (no parsing) or a `ShaderReflection` (for shaders compiled at run time). ```csharp using OneShader.Layouts; // Built with the project: one call per target. PipelineLayout pipeline = Shaders.Sprite.Pipeline; // vertex + pixel entry points of the file var vulkan = pipeline.ToVulkan(); // SetLayouts[i].Bindings, PushConstantRanges foreach (var set in vulkan.SetLayouts) foreach (var b in set.Bindings) Console.WriteLine($"set {set.Set} binding {b.Binding}: {b.DescriptorType} x{b.DescriptorCount} stages {b.StageFlags}"); var vertex = Shaders.Sprite.VSMain.Layout.Vertex!.ToVulkan(VertexStepMode.Instance); // VkVertexInputBinding/AttributeDescription data // Any combination, and conflicts come back as values: if (!PipelineLayout.TryCreate([Quad.VSMain.Layout, Quad.PSBlur.Layout], out var blur, out var error)) Console.WriteLine(error); // Compiled at run time: EntryLayout.TryCreate(result.Reflection!, "PSMain", out var ps, out _); ``` Specialization constants (`[[vk::constant_id(N)]] const T Name = default;`) are pipeline-time inputs on every target. `SpecializationValues.For(entryLayout)` holds one value per constant the entry point reads (defaults until you `Set` them) and converts to `VkSpecializationInfo` contents (`ToVulkan()`), Metal function-constant values (`ToMetal()`) and WebGPU's `constants` record (`ToWebGpu()`); the generated `Specialization` struct does the same by field. To fix values at compile time instead, pass `CompileOptions.SpecializationOverrides`. `PipelineLayout.TryCreate` merges entry points that run together: bindings at the same set and binding must agree on kind, count and type and get the union of their stages; push-constant blocks of the same name and size are shared. It returns a `LayoutError` (`ResourceConflict`, `DuplicateStage`, ...) with the message to show, never an exception. The three conversions: - `ToVulkan(VulkanLayoutOptions?)`: a `DescriptorSetLayoutInfo` per set from 0 (unused sets are empty, as Vulkan needs them), and a `PushConstantRangeInfo` per block, using the numeric values of `VkDescriptorType`, `VkShaderStageFlagBits` and `VkFormat`. `DynamicBuffers` names buffers to declare as dynamic-offset descriptors. - `TryToMetal(...)`: for each stage, the `[[buffer(n)]]`, `[[texture(n)]]` and `[[sampler(n)]]` index of every resource (an array of N takes N indices), the buffer index to pass to `setVertexBytes`/`setFragmentBytes`/`setBytes`, and `MetalStageLayout.TryFind(set, binding)` to go from a Vulkan-style slot to the index. `VertexLayout.ToMetal(bufferIndex, step)` gives the `MTLVertexDescriptor` data (default buffer index 30; pass the generated `MslVertexBufferIndex`). Fails with `MissingTarget` if `Msl` was not built. - `TryToWebGpu(..., WebGpuLayoutOptions?)`: `GPUBindGroupLayoutEntry` data (`ToWebGpuString()` gives the API strings), with the push-constant block as a uniform buffer entry at the group and binding the compiler reserved, and `VertexLayout.ToWebGpu(step)` for `GPUVertexBufferLayout`. What WebGPU cannot express (resource arrays, typed buffers, 1D array textures, storage textures without `[[vk::image_format]]`) is an `Unsupported` error that says what to change. Reflection does not record whether a texture is a depth texture sampled with a comparison sampler, or how a storage texture is accessed, so those come from `WebGpuLayoutOptions` (`DepthTextures`, `StorageAccess`, ...); defaults are `float` textures and `write-only` storage textures. One layout for many pipelines: ```csharp // A material system shares one layout across ten forward entry points and three post-effect pixel shaders (several per stage). if (!PipelineLayout.TryCreateUnion(allEntries, out var shared, out var error)) Console.WriteLine(error); // "Set 1 binding 3 is 'Albedo' (SampledTexture, ...) in entry point 'PsTex' and 'Params' (ConstantBuffer, ...) in entry point 'PsTint'. ..." // A bind group made once has to fit pipelines that use different stages: var vk = shared.ToVulkan(new VulkanLayoutOptions { WidenSets = new Dictionary { [0] = ShaderStages.Vertex | ShaderStages.Pixel }, DynamicBindings = [(0u, 1u)], // by slot; DynamicBuffers = ["Frame"] does it by name }); var web = shared.TryToWebGpu(out var wl, out _, new WebGpuLayoutOptions { WidenGroups = ..., DynamicBuffers = ["Frame"] }); // The engine stores the colour as unorm8x4 at its own offset in a 48-byte vertex: var overrides = new VertexLayoutOverrides { Stride = 48, Attributes = new Dictionary { [4] = new(VertexFormat.Unorm8x4, 44) } }; if (Quad.VSMain.Layout.TryWithVertexOverrides(overrides, out var vs, out error)) // then use vs in TryCreate/TryCreateUnion Console.WriteLine(vs!.Vertex!.ToMetal().Attributes[4]); // also ToVulkan() and ToWebGpu() ``` `TryCreateUnion` merges resources at one set and binding that agree in kind, count, type, texel type and storage format (stages are united); a disagreement names both entry points and the slot. Push-constant blocks of the same name and size are shared; blocks of different structs merge into one range covering all their stages, sized to the largest, when the members that overlap in bytes have the same type at the same offset (this needs the layout built by `EntryLayout.TryCreate`, which carries `PushConstantType`; generated layouts can only merge blocks of the same name and size). Entry points of different files number Metal resources independently, so `TryToMetal` reports a slot or push block on which two entry points of one stage disagree (Vulkan and WebGPU are unaffected, apart from push-constant slots in the WGSL). Options that shape the output without changing what the shader says: `WidenSets`/`WidenBindings`/`WidenPushConstants` (Vulkan) and `WidenGroups`/`WidenBindings`/`WidenPushConstants` (WebGPU) OR extra stages into the visibility; `DynamicBuffers` (by name), `DynamicBindings` (by set and binding) and `ResourceLayout.DynamicOffset` (on a layout you built) make a constant or storage buffer `*_BUFFER_DYNAMIC` / `hasDynamicOffset` (`DescriptorSetLayoutBindingInfo.IsDynamic`, `BindGroupLayoutEntryInfo.HasDynamicOffset`). `TryWithVertexOverrides` accepts a format of the same class as the shader's input (float, unorm, snorm and half formats for float inputs, Uint formats for `uint`, Sint formats for `int`); a different component count is fine, because Vulkan, Metal and WebGPU all drop or fill components. The attribute must lie inside the stride, and its offset must be a multiple of min(4, format size). Otherwise `LayoutErrorKind.InvalidVertexInput`. Creating layouts allocates; do it when a shader loads, not per frame. ## What works - Vertex, pixel and compute shaders in modern HLSL (HLSL 2021): structs, functions, constant buffers, textures and samplers, structured and byte-address buffers, storage images, groupshared memory, barriers and atomics, and the preprocessor (`#include`, macros, `#if`). - SPIR-V for Vulkan 1.2, validated with `spirv-val`, and checked by running each shader in its test corpus on a GPU against the same shader compiled by DXC. - Buffer layouts and descriptor bindings follow the conventions DXC uses for Vulkan. - Metal Shading Language (`CompileTarget.Msl`, Metal 2.4+) and WGSL (`CompileTarget.Wgsl`) with the same buffer byte layouts as the SPIR-V, checked with Apple's `metal` compiler and `naga`, and by running the test corpus on Metal and in a browser on WebGPU against the DXC reference. - Push constants (`[[vk::push_constant]]` on a struct global or `ConstantBuffer`): a `PushConstant` block in SPIR-V with DXC's offsets, a `constant T&` argument in MSL, a uniform buffer in WGSL; `ShaderReflection.PushConstants` lists the layout, the stages and where each target takes the bytes. ## Known limits - SPIR-V, Metal (MSL) and WebGPU (WGSL) output. Limits are listed in `docs/msl-backend.md` and `docs/wgsl-backend.md`. - Unbounded resource arrays and `WaveMatch`/`WaveMultiPrefix*` are not supported. Wave and quad intrinsics work; a shader that uses them needs subgroup support on the device (`EntryPointInfo.SubgroupOperations`). ## License MIT OR Apache-2.0, at your option. Copyright © Recursive AI, LLC. # OneShader.Tool **`oneshader`: the OneShader HLSL compiler on the command line.** It compiles a vertex, pixel or compute shader to SPIR-V (Vulkan 1.2+), Metal Shading Language or WGSL, and takes DXC's core flags, so a DXC command line mostly works unchanged. Use it from CMake, Make, scripts or any build system that is not MSBuild. (In an MSBuild project, [RecursiveAI.OneShader.Build](https://www.nuget.org/packages/RecursiveAI.OneShader.Build) compiles shaders during `dotnet build` and generates C# bindings.) ## Install ```sh dotnet tool install -g RecursiveAI.OneShader.Tool oneshader --version ``` Needs the .NET 10 runtime. For a single native binary with no runtime and the fastest start-up (about 2.4 times faster per process than DXC), publish it with NativeAOT from this repository: ```sh dotnet publish src/OneShader.Tool -c Release -r linux-x64 -o out # out/oneshader ``` ## Use ```sh oneshader -T ps_6_0 -E PSMain -spirv sprite.hlsl -Fo sprite.ps.spv oneshader --target msl --target wgsl --all-entries sprite.hlsl -Fo build/sprite.out oneshader -T cs_6_0 -E CSMain -D TILE=16 -I include --reflect histogram.json --depfile histogram.d histogram.hlsl -Fo histogram.spv ``` Diagnostics go to standard error as `file(line,col): error OSH0102: message`, the form IDEs and CMake parse. Exit code: **0** success, **1** the shader did not compile (or an output could not be written), **2** the command line is wrong. Warnings print and do not change the exit code unless `-WX` is given. ### DXC-compatible options | Option | Meaning | |---|---| | `-T ` | `vs_6_x`, `ps_6_x`, `cs_6_x` (also `_5_x`); the stage comes from the profile, the model number is accepted and otherwise unused. Geometry, tessellation, mesh and library profiles are rejected. | | `-E ` | Entry point; with `-T` and no `-E` it is `main`, as in DXC. | | `-D NAME[=VALUE]` | Macro; without a value it is `1`. | | `-I ` | `#include` search directory. | | `-Fo ` | Output file. Several outputs (entry points × targets) are written beside it: `name.` when there are several entry points, and `.spv`/`.metal`/`.wgsl` when there are several targets. Without `-Fo`, a single output goes to standard output. | | `-spirv` | Emit SPIR-V. This is also the default when no target is named (DXC's default is DXIL, which is not implemented). | | `-O0`, `-Od` / `-O1`..`-O3` | Optimizer off / on (default on). | | `-Zpr` / `-Zpc` | Row-major / column-major default matrix packing. | | `-HV 2021` | Accepted (`202x` too); other versions are rejected. | | `-WX` | Warnings are errors. | | `-enable-16bit-types` | `half` and `min16*` are 16-bit. | | `-fspv-target-env=vulkan1.2` | Accepted (`vulkan1.3` too); the output is valid for Vulkan 1.2 and newer. | | `-nologo` | Accepted, ignored. | Values may follow the flag after a space or directly (`-T ps_6_0`, `-Tps_6_0`). Every other `-fvk-*` and `-fspv-*` option, and DXC options such as `-Zi` or `-Fh`, is rejected with a message naming it. oneshader always uses DXC's default Vulkan layout and binding rules; use `[[vk::binding]]`, `[[vk::location]]` and `[[vk::offset]]` in the shader. ### oneshader options | Option | Meaning | |---|---| | `--target spirv\|msl\|wgsl` | Target; repeat for several. | | `--all-entries` | Compile every `[shader("vertex"\|"pixel"\|"compute")]` entry point. This is also what happens without `-T`. Cannot be combined with `-T`/`-E`. | | `--reflect ` | Write reflection as JSON (schema below). | | `--depfile ` | Write a Make-style dependency file, `outputs: files read`, with `$`, `#` and spaces escaped. Written on success. | | `--no-names` | Leave source names out of the output. | | `--version`, `--help` | | ## Reflection JSON, schema version 1 The same data as the library's `ShaderReflection`, keys in camelCase, enums by name, output in a fixed order, `\n` line ends and no machine paths, so the same compile gives the same bytes. Changes that break readers raise `schemaVersion`. ``` { "schemaVersion": 1, "source": "", "targets": ["SpirV", ...], "resources": [ { "name", "kind", "hlslType", "resourceType", "count", "stages": ["Vertex"|"Pixel"|"Compute"], "slots": [ { "target", "set", "binding" } ], "counterSlots": [ ... ], "layout": } ], "pushConstants": [ { "name", "stages", "slots", "layout": } ], "specConstants": [ { "name", "id", "type": "Bool"|"Int"|"UInt"|"Float", "defaultBits", "slots": [ { "target", "index", "name" } ] } ] (only when the shader has some), "entryPoints": [ { "name", "stage", "workgroupSize": [x, y, z] (compute), "vertexStride" (vertex), "inputs": [ ], "outputs": [ ], "usedResources": [ names ], "pushConstant": name (optional), "usedSpecConstants": [ names ] (optional) } ] } = { "name", "semantic", "semanticIndex", "location" (optional), "builtIn", "scalar", "components", "format" (optional), "offset", "interpolation" } = { "kind": "Scalar|Vector|Matrix|Array|Struct", "size", "alignment", and by kind: Scalar: "scalar" | Vector: "scalar", "components" | Matrix: "scalar", "rows", "columns", "matrixStride", "packing" Array: "length", "stride", "element": | Struct: "name", "members": [ { "name", "offset", "type": } ] } ``` Sizes, offsets and strides are the byte layout of the buffer, equal for every target. Licensed under MIT OR Apache-2.0. # Generated bindings The shape below is what the build package generates. Changes to it are listed in the changelog. ## Adding shaders to a project Reference the `RecursiveAI.OneShader.Build` package. Every `*.hlsl` file in the project becomes a `OneShader` item and is compiled during `dotnet build` (turn the default glob off with `false` and list items yourself). ```xml ``` The package brings `RecursiveAI.OneShader` (the support types) with it. Build with `dotnet build`: the task is a .NET task, so Visual Studio's .NET Framework MSBuild cannot run it. To use a local build instead of NuGet, `tools/pack-local.sh` packs both packages into a local folder (`.feed` by default); `samples/` has projects that restore from it, and one (`samples/LocalSource`) that references the repository's projects directly. Only projects that contain shaders reference the build package; it does not flow to their consumers. Entry points are found from `[shader("vertex")]`, `[shader("pixel")]` and `[shader("compute")]` attributes. Item metadata can list them explicitly instead. Item metadata, all optional: `EntryPoints="VSMain:Vertex;PSMain:Pixel"`, `Defines="A;B=2"`, `IncludeDirectories="dir;dir"` (searched after the including file's folder), `Namespace`, `ClassName`, `Targets`, and for variants `Axes` and `MaxVariants` (see Variants below). Project properties: `OneShaderNamespace` (default `$(RootNamespace).Shaders`), `OneShaderDefines`, `OneShaderIncludeDirectories`, `OneShaderMaxVariants`. Listing a file as a `OneShader` item (to give it metadata) is fine: the default glob skips files that are already listed. A shader is compiled again only when it, a file it includes, the options or the compiler changed. Errors and warnings are ordinary MSBuild diagnostics on the HLSL line: ``` Shaders/Sprite.hlsl(12,17): error OSH0300: 'Atlass' is not declared. Did you mean 'Atlas'? ``` ## What is generated One `static partial class` per HLSL file, named after the file (see Names below), in the namespace `.Shaders` (overridable). Nothing uses reflection or file I/O; output is deterministic. - **Constant buffers**: one blittable struct per `cbuffer` / `ConstantBuffer`, named after it, with `[StructLayout(LayoutKind.Explicit)]`, a `uint Size` constant, and field offsets equal to the offsets in the compiled shader. Float vectors and matrices use `System.Numerics` types. - **Bindings**: a nested `Bindings` class with, per resource, `const uint Set`, `Binding`, `Count` and `const ResourceKind Kind`, plus `All`, a span of `ResourceBinding` that also says which stages use each resource. See Resource arrays and Append and consume buffers below. - **Entry points**: one nested class per entry point, named after the function, with `Stage`, `EntryPoint`, and `SpirV` (a `ReadOnlySpan` over data embedded in the assembly). - Vertex entry points add a blittable `Vertex` struct, `uint Stride`, and `Inputs`, a span of `VertexAttribute` (location, `VertexFormat`, offset). - Compute entry points add `ThreadGroupSizeX`, `ThreadGroupSizeY` and `ThreadGroupSizeZ`. - With the Msl or Wgsl target selected, entry points also have `Msl` and `Wgsl` members. - **Reflection**: the file's class has `static ShaderReflection Reflection`, the reflection the rest of the class was generated from (resources with slots and layouts, entry points, vertex inputs, thread-group sizes). It is stored as text in the assembly (`ReadOnlySpan ReflectionData`, about 1 to 4 KB per shader) and parsed on the first read, so a program that never reads it carries no parser after trimming. Hot reload compares it with a fresh compile; see "Hot reload" in the `RecursiveAI.OneShader` package README. - All counts, sizes, strides, offsets, locations and thread-group sizes are `uint`, matching the Vulkan API. ### Push constants A struct-typed global or `ConstantBuffer` marked `[[vk::push_constant]]` is a push-constant block, not a descriptor: ```hlsl struct PerDraw { float4x4 Model; float4 Tint; uint Index; }; [[vk::push_constant]] ConstantBuffer Draw; // or: [[vk::push_constant]] PerDraw Draw; ``` The generated class gets a blittable `PushConstants` struct (with `Size` and the SPIR-V member offsets) and `Bindings.PushConstants` with the delivery information; the block is not in `Bindings.All`. ```csharp public static class PushConstants { public const uint Size = 96u; public const ShaderStages Stages = ShaderStages.Vertex | ShaderStages.Pixel; public const uint MslIndex = 0u; // when Msl is selected public const uint WgslGroup = 0u, WgslBinding = 2u; } // when Wgsl is selected ``` - **Vulkan:** `vkCmdPushConstants(cmd, layout, (VkShaderStageFlags)Bindings.PushConstants.Stages, 0, Size, &data)`; the pipeline layout's range is offset 0, `Size` bytes, those stages. Vulkan only guarantees 128 bytes; a larger block is warning OSH0334. - **Metal:** `setVertexBytes` / `setFragmentBytes` (or `setBytes` for compute) with the struct's bytes at `MslIndex`. - **WebGPU:** a uniform buffer at `@group(WgslGroup) @binding(WgslBinding)`, at least `Size` rounded up to 16 bytes long. The binding is the one after the highest binding of the highest bind group the shader uses, so it never collides. - Members follow the vector-relaxed std430 rules DXC uses for push constants (a `float3` may share a 16-byte row with the next scalar), not the constant-buffer rules; the same bytes are valid on every target. - A file may declare several blocks as long as no entry point reads more than one (OSH0333). With one block the struct and the `Bindings` class are named `PushConstants`; with several, each is named after its variable, `PushConstants`. ### Resource arrays `Texture2D Tex[4]`, `SamplerState Smp[2]`, arrays of structured buffers and `ConstantBuffer Cb[2]` are bounded arrays of resources. An array is **one** descriptor binding with a descriptor count, as in DXC's SPIR-V (the next unnumbered resource gets the next binding number, not `Binding + Count`). ```csharp public static class Tex { public const uint Set = 0, Binding = 0, Count = 4; public const ResourceKind Kind = ResourceKind.SampledTexture; } // All contains: new ResourceBinding("Tex", 0, 0, 4, ResourceKind.SampledTexture, ShaderStages.Compute) ``` `Count` is 1 for a single resource. The shader may index an array with a constant or a run-time value (pass the elements to helper functions too); `NonUniformResourceIndex(i)` marks an index that differs between invocations. Run-time indices need the Vulkan device features `shaderSampledImageArrayDynamicIndexing`, `shaderStorageBufferArrayDynamicIndexing`, `shaderUniformBufferArrayDynamicIndexing` and `shaderStorageImageArrayDynamicIndexing` (the compiled shader declares the matching capability for the kinds it indexes), and `shaderSampledImageArrayNonUniformIndexing` / `shaderStorageBufferArrayNonUniformIndexing` for non-uniform indices. Unbounded arrays (`Tex[]`) are an error (`OSH0350`); multi-dimensional resource arrays are an error too. ### Append and consume buffers `AppendStructuredBuffer` and `ConsumeStructuredBuffer` take two descriptors: the element buffer and a hidden 4-byte counter buffer in the **same set**. The counter holds the number of elements: `Append` adds one and stores at the old value, `Consume` subtracts one and reads the element below the old value. `GetDimensions` reports the element buffer's capacity and stride, not the counter. The application owns both buffers: bind a 4-byte (or larger) storage buffer at the counter binding, zero it to start an append list empty, and write the element count into it to start a consume list full. The counter binding is `[[vk::counter_binding(n)]]` if given; otherwise it is the next free binding of the buffer's set after the resources declared so far (the same numbering DXC uses), so `AppendStructuredBuffer A : register(u0)` next to `B : register(u1)` puts `A`'s counter at 2 and `B`'s at 3 if those are free. ```csharp public static class Hits { public const uint Set = 0, Binding = 2, Count = 1; public const ResourceKind Kind = ResourceKind.StorageBuffer; public const uint CounterBinding = 3; // same Set } // All contains the buffer and, right after it, its counter: // new ResourceBinding("Hits", 0, 2, 1, ResourceKind.StorageBuffer, ShaderStages.Compute) // new ResourceBinding("Hits.counter", 0, 3, 1, ResourceKind.StorageBuffer, ShaderStages.Compute) ``` A resource without a counter has no `CounterBinding` constant. Arrays of append or consume buffers are an error. `RWStructuredBuffer.IncrementCounter` and `DecrementCounter` are not supported yet. The names `Reflection` and `ReflectionData` are taken in the file's class, so a constant buffer, struct or entry point with one of those names is a name-clash error. Member types in generated structs: - `float`, `int`, `uint` and the other scalars map to the C# type of the same width; `bool` is a `uint` (HLSL stores it as a 32-bit word). - `float2/3/4` are `Vector2/3/4`. Vectors of other element types have no `System.Numerics` type and become one field per component (`Offset_X`, `Offset_Y`). - Arrays become one member per element (`Weights_0`, `Weights_1`), so the struct needs no unsafe code. - Nested HLSL structs become nested C# structs with the same rules. If one HLSL struct appears under two packings (a constant buffer and a structured buffer), the second gets a `_1` suffix. - Matrices are `Matrix4x4` members, for every size up to 4x4 (elements outside the HLSL shape stay zero). A row-major 4x4 is a plain field; the other shapes are properties over private storage that does the arranging, so element (i, j) is `M[i][j]` either way. - Each `StructuredBuffer` / `RWStructuredBuffer` whose `T` is a struct also gets a C# struct, named after `T`, with `Size` equal to the element stride. - Vertex fields follow the format: `Float32x2` is a `Vector2`; packed formats (`Unorm8x4` and the like) are a `uint` or `ulong` holding the raw bits; integer and half vectors are one field per component. The support types (`ShaderStage`, `ResourceKind`, `VertexFormat`, `VertexAttribute`, `ResourceBinding`) live in `RecursiveAI.OneShader`, which the build package brings in, so a rendering layer in another assembly can accept them. When only these types are used, trimming removes the compiler from a NativeAOT app. ## Specialization constants A file that declares `[[vk::constant_id(N)]] const T Name = default;` constants gets, next to the bindings: ```csharp public static class SpecializationConstants { public static class Rounds { public const uint Id = 1u; public const int Default = 4; } /* ... */ } public struct Specialization { public uint UseOffset; public int Rounds; public uint Mask; public float Scale; // 4 bytes each, declaration order; a bool is 0 or 1 public static Specialization Defaults { get; } public SpecializationInfo ToVulkan(ReadOnlySpan ids = default); // map entries + data for VkSpecializationInfo public IReadOnlyList ToMetal(ReadOnlySpan ids = default); public IReadOnlyDictionary ToWebGpu(ReadOnlySpan ids = default); } // keys: decimal ids public static class CSMain { public static ReadOnlySpan SpecConstantIds { get; } /* ids this entry point reads */ } ``` Pass the entry point's `SpecConstantIds` to the `To...` methods so a pipeline gets only the constants its entry point reads (WebGPU rejects others; Vulkan and Metal ignore them). `CSMain.Layout.SpecConstants` lists them with names, types and defaults, and `SpecializationValues.For(CSMain.Layout)` is the same data driven by names (`Set("Rounds", 7)`) instead of struct fields. Entry points that read no constant have no `SpecConstantIds`. Names that would clash with a member of `Specialization` (`Size`, `Count`, `Defaults`, ...) are reported by the generator. Metal hosts: a function with function constants must be created with `makeFunction(name:constantValues:)` even to get the defaults (see the MSL backend notes). `CompileOptions.SpecializationOverrides` (runtime compiler API) fixes constants at compile time instead. In a variant file the same `SpecializationConstants` and `Specialization` types are generated inside each variant (`Variants..Specialization`), because the set of constants can differ per variant; the aggregate `Entry.SpecConstantIds` returns the ids the looked-up variant reads. ## Variants A `OneShader` item with `Axes="ALPHA_TEST;LIGHTING=NONE|LAMBERT|PBR"`, or `OneShaderVariant` items naming the shader, is compiled once per define set (see the build package readme for the MSBuild side). The file's class then has this shape instead of the single-variant members `Bindings`, `Reflection`, `Pipeline` and the entry-point classes at file level (those move into each variant): | Member | Meaning | |---|---| | `Frame`, `PushConstants`, `Vertex`, other structs | generated once when every variant agrees on the layout | | `Value` | `enum` per choice axis (`LightingValue { None, Lambert, Pbr }`); flags are `bool`; named sets use `VariantName` | | `Key` | `readonly struct`, one property per axis, `Index` (-1 for an enum value out of range), equality | | `Variant` | `Index`, `Name`, `Defines`, `Key`, `Bindings`, `Reflection`, `Pipeline` (when any variant has one) and one `Entry` property per entry point | | `Entry` | `Stage`, `EntryPoint`, `SpirV`, `Msl`, `MslFunctionName`, `MslVertexBufferIndex`, `Wgsl`, `WgslEntryPoint`, `Stride`, `Inputs`, `ThreadGroupSizeX/Y/Z`, `SubgroupOperations`, `RequiredFeatures`, `Layout` (each only when some variant has it) | | `Variants` | `Count`, `All`, `Get(Key)`, `Get(int)`, `TryGet(Key, out Variant)`, and one static class per variant (`Variants.AlphaTestOn_LightingPbr`) shaped like a single-variant file's class | Variant order is the cartesian product with the first axis changing slowest; `Key.Index` is the position in `Variants.All`. A flag axis contributes `On`/`Off` to the variant class name, a choice contributes ``, joined by `_`; all-capitals macro names and values are PascalCased (`ALPHA_TEST` gives `AlphaTest`). Lookup is an array index and the blobs are returned through switches over constants: no allocation, no reflection, trim- and AOT-clean. **When variants disagree about a struct.** Two variants may lay out a `cbuffer` (or push block, or vertex input) differently, for example because a member sits under `#if`. Then no shared type is generated for that name; each variant's class declares its own (`Variants..Frame`, `Variants..VSMain.Vertex`), matching that variant's reflection. Resource slots that differ between variants are not an error: read them from `Variant.Bindings` and `Variant.Pipeline`. ## Pipeline layouts Each entry-point class has `Layout`, an `OneShader.Layouts.EntryLayout` with the resources that entry point uses (with its own stage only), its push-constant block, its vertex input and its thread-group size, with the slot on every selected target. The file's class has `Pipeline` when the file has exactly one vertex and one pixel entry point, or exactly one compute entry point, and no cbuffer, struct or entry point of that name. For other combinations, merge entry points yourself with `PipelineLayout.TryCreate`, or with `PipelineLayout.TryCreateUnion` when one layout must serve many entry points, several per stage. ```csharp public static class VSMain { public static EntryLayout Layout { get; } /* ... */ } public static PipelineLayout Pipeline { get; } // Create(VSMain.Layout, PSMain.Layout) ``` Both are built on first use from constant arrays (no reflection parsing) and cached. `Pipeline.ToVulkan()`, `TryToMetal()` and `TryToWebGpu()` convert to the target's description; see the `RecursiveAI.OneShader` readme. ## Names C# warns about type names made only of lower-case ASCII letters (CS8981), which is an error in projects that treat warnings as errors. So every generated **type** name (the file's class, constant buffers, structs, the classes under `Bindings`, entry-point classes) follows one rule: - A name with an upper-case letter keeps its spelling: `Quad`, `Tile`, `UiRect`, `Frame`. Characters that cannot be in an identifier (`-`, `.`, space) become `_`: `UI-rect` gives `UI_rect`. - Any other name is turned into PascalCase: split at `_` and at characters that cannot be in an identifier, upper-case the first letter of each part, join. `sprite` gives `Sprite`, `msdf_text` gives `MsdfText`, `ui-rect` gives `UiRect`, `vs_main` gives `VsMain`. - A name that starts with a digit gets a `_` in front: `2d_sprite` gives `_2dSprite`. - The `ClassName` item metadata is used exactly as written. Strings are never changed: `EntryPoint` is still `"vs_main"`, and `Bindings.All` lists the HLSL resource names. Members (struct fields, properties, constants) keep their HLSL spelling; a C# keyword is escaped with `@` (`@class`). Two names that map to the same C# name, such as `ui-rect` and `ui_rect`, are a build error that says to rename one or set `ClassName`. ## 16-bit and 64-bit members A buffer member of a sized type is generated at its DXC offset with the matching C# type: `half` is `System.Half`, `int16_t` `short`, `uint16_t` `ushort`, `int64_t` `long`, `uint64_t` `ulong`, `double` `double` (min-precision types without 16-bit types are `float`/`int`/`uint`). Vectors of these have no `System.Numerics` form, so they are one field per component: `half3 b` becomes `b_X`, `b_Y`, `b_Z` (each at its own offset). A matrix of a 16-bit or 64-bit element is an error that says to split it (only float matrices up to 4x4 are generated). Each entry class gets `RequiredFeatures` (a `DeviceFeatures` constant) beside `SubgroupOperations`, and its `Layout` carries it; enable the Vulkan features it names (or request `shader-f16` on WebGPU) when creating the device. Turn the types on with `true` or `Enable16BitTypes="true"` on a `OneShader` item. ## Layout rules - One byte layout per buffer on every target: HLSL packing as used for Vulkan (relaxed std140 for constant buffers, relaxed std430 for structured buffers); `packoffset` and `[[vk::offset]]` are honoured. The same C# struct is valid for every backend. - Matrices are `column_major` by default, as in HLSL; `row_major` and `#pragma pack_matrix` are honoured. The generated C# member is arranged so that element (i, j) in C# is `M[i][j]` in HLSL whichever packing the shader uses. - Vulkan descriptor slots: `[[vk::binding(binding, set)]]` wins. Otherwise set is the register space and binding is the register number. Two resources on the same slot (for example `t0` and `s0`) are a build error that says to add `[[vk::binding]]` or renumber. Resources with neither get set 0 and the lowest binding number not taken by an explicit slot, in declaration order (what DXC does). Plain globals outside any `cbuffer` collect in an implicit constant buffer named `Globals` (`_Globals` if that name is taken). - Vertex input locations follow declaration order unless `[[vk::location]]` is given. Formats follow the shader type; a packed format is declared on the input field with `[[oneshader::format("unorm8x4")]]` (any `VertexFormat` name, case-insensitive; the component count and integer/float kind must match the declared type, e.g. `float4` for `unorm8x4`). The shader sees the converted `float4`. DXC compiles the same file with a warning ("unknown attribute 'format' ignored") and no error. Offsets are packed in declaration order, so the vertex struct follows the format as described under "Vertex fields follow the format". ## Example ```hlsl // Shaders/Sprite.hlsl cbuffer Frame : register(b0) { float4x4 ViewProjection; }; Texture2D Atlas : register(t1); SamplerState AtlasSampler : register(s2); struct VSIn { float2 Position : POSITION; float2 Uv : TEXCOORD0; float4 Color : COLOR0; }; struct VSOut { float4 Position : SV_Position; float2 Uv : TEXCOORD0; float4 Color : COLOR0; }; [shader("vertex")] VSOut VSMain(VSIn i) { VSOut o; o.Position = mul(ViewProjection, float4(i.Position, 0, 1)); o.Uv = i.Uv; o.Color = i.Color; return o; } [shader("pixel")] float4 PSMain(VSOut i) : SV_Target { return Atlas.Sample(AtlasSampler, i.Uv) * i.Color; } ``` generates, in outline: ```csharp namespace MyGame.Shaders; public static partial class Sprite { [StructLayout(LayoutKind.Explicit, Size = 64)] public struct Frame { public const uint Size = 64; /* ViewProjection at offset 0 */ } public static class Bindings { public static class Frame { public const uint Set = 0, Binding = 0, Count = 1; public const ResourceKind Kind = ResourceKind.ConstantBuffer; } public static class Atlas { public const uint Set = 0, Binding = 1, Count = 1; public const ResourceKind Kind = ResourceKind.SampledTexture; } public static class AtlasSampler { public const uint Set = 0, Binding = 2, Count = 1; public const ResourceKind Kind = ResourceKind.Sampler; } public static ReadOnlySpan All => /* ... */; } public static class VSMain { public const ShaderStage Stage = ShaderStage.Vertex; public const string EntryPoint = "VSMain"; public static ReadOnlySpan SpirV => /* embedded */; [StructLayout(LayoutKind.Explicit, Size = 32)] public struct Vertex { /* Position @0 Vector2, Uv @8 Vector2, Color @16 Vector4 */ } public const uint Stride = 32; public static ReadOnlySpan Inputs => /* (0, Float32x2, 0), (1, Float32x2, 8), (2, Float32x4, 16) */; } public static class PSMain { public const ShaderStage Stage = ShaderStage.Pixel; public const string EntryPoint = "PSMain"; public static ReadOnlySpan SpirV => /* embedded */; } } ``` # Reflection file format The build task writes one reflection file per shader; the source generator reads it without loading the compiler. The model is `OneShader.Reflection.ShaderReflection`; `ReflectionFormat.Write` and `TryRead` implement this format. Version 1. ## Rules - UTF-8, no byte-order mark, `\n` line ends, one record per line, ending with a newline. - First line: `oneshader-reflection 1`. A reader must reject other versions. Additive changes (new records or new attributes) keep version 1; readers ignore unknown records and attributes. - A record is `keyword key=value key=value ...`. Values are bare tokens (no spaces) or double-quoted strings with `\\`, `\"`, `\n`, `\r`, `\t` and `\uXXXX` escapes. Numbers are invariant-culture decimals; all counts, sizes, strides, offsets, locations and workgroup sizes are unsigned. Enum values use the member names of the C# enums. Lists are comma-separated; an empty list is `-`. - Nesting is by indentation: two spaces per level. - Order is the model's order (declaration order). Equal models give identical bytes. ## Records ``` source name="Shaders/Sprite.hlsl" targets list=SpirV,Wgsl resource name="Frame" kind=ConstantBuffer count=1 stages=Vertex,Pixel rtype=ConstantBuffer hlsl="" slot target=SpirV set=0 binding=0 layout type kind=Struct scalar=Bool rows=0 cols=0 name="Frame" size=128 align=16 length=0 stride=0 mstride=0 packing=ColumnMajor member name="ViewProjection" offset=0 type kind=Matrix ... mstride=16 packing=ColumnMajor member name="Weights" offset=80 type kind=Array ... length=3 stride=16 element type kind=Scalar scalar=Float ... entry name="VSMain" stage=Vertex workgroup=0,0,0 stride=32 input name="Position" semantic="POSITION" index=0 location=0 builtin=None scalar=Float components=2 format=Float32x2 offset=0 interp=Default output name="Position" semantic="SV_Position" index=0 location=- builtin=Position scalar=Float components=4 format=None offset=0 interp=Default uses name="Frame" ``` - `resource`: an optional `format` attribute holds the `[[vk::image_format]]` name of a read-write texture; `kind` is `ResourceKind`; `stages` a list of `Vertex`, `Pixel`, `Compute` (`-` for none); `rtype` is `ResourceTypeKind`; `hlsl` the type as written. Child `slot` records give one slot per target (`target` is `SpirV`, `Msl` or `Wgsl`; Metal uses the argument index as `binding` and 0 as `set`). `count` is the array size (1 for a single resource); an array occupies one `slot`. A `layout` child holds a `type` tree. - `counter` (child of `resource`, append and consume buffers only): `target`, `set`, `binding` of the hidden counter buffer, one record per target like `slot`. Readers that do not know it ignore it. - `push`: a push-constant block (not a `resource`): `name`, `stages`, a `layout` child with a `type` tree, and `slot` children for the targets that deliver it as a bound buffer: `Msl` (`binding` is the `[[buffer(n)]]` index, `set` 0) and `Wgsl` (`set` is the bind group, `binding` the binding). SPIR-V has no slot. An `entry` names the block it reads with a `pushuse name=...` child. - `features` (child of `entry`, written only when non-zero): `bits`, the `DeviceFeatures` flags the entry point needs for 16-bit and 64-bit types (1 Float16, 2 Int16, 4 Float64, 8 Int64, 16 StorageBuffer16, 32 UniformBuffer16, 64 PushConstant16, 128 InputOutput16, 256 Int64Atomics). - `spec`: a specialization constant: `name`, `id` (the `constant_id`), `scalar` (`Bool`, `Int`, `UInt` or `Float`) and `default` (the 32-bit pattern, as a number). Child `slot` records (`target`, `index`, `name`) give the SPIR-V `SpecId`, the Metal function-constant index and variable, and the WGSL `@id` and override name per target. Written only for files that have constants; it follows the `push` records. An `entry` lists the constants it reads with `specuse name=...` children. - `type`: `kind` is `Scalar`, `Vector`, `Matrix`, `Array` or `Struct`. `rows` is the vector component count or matrix rows; `cols` matrix columns; `size`, `align` in bytes; `length`/`stride` for arrays (length 0 = unsized); `mstride`/`packing` for matrices (the stride is between stored columns for `ColumnMajor`, rows for `RowMajor`). An `Array` has one `element` child holding a `type`; a `Struct` has `member` children (`name`, `offset`) each holding one `type`. - `entry`: `workgroup` is `x,y,z` (zeros unless compute); `stride` the interleaved vertex size (vertex stage only). `input`/`output` are stage variables; `location=-` means built-in; `format` is a `VertexFormat` name for vertex inputs and `None` elsewhere. `uses` lists resource names the entry point touches. Layouts are the single buffer layout shared by all targets (see `BufferLayout`). # The SPIR-V backend `SpirvBackend.Emit(IrModule, entryPointName, SpirvOptions?, DiagnosticBag)` turns one entry point of an IR module, and everything it reaches, into a SPIR-V 1.5 binary for the Vulkan 1.2 environment (Shader capability, logical addressing, GLSL450 memory model, GLSL.std.450 for math). It returns `null` and reports `OSH50xx` errors when something cannot be expressed; it does not throw for unsupported input. The same module and options always give the same bytes. Code is in `src/OneShader/SpirV/`. ## Mapping decisions - **Matrices are transposed.** An HLSL `floatRxC` is a SPIR-V matrix with R columns of C components, so `m[i]` is column `i`, and `mul(a, b)` is `OpMatrixTimesMatrix(b, a)`, `mul(v, m)` is `OpMatrixTimesVector(m, v)` and `mul(m, v)` is `OpVectorTimesMatrix(v, m)`. In buffers, HLSL `column_major` becomes `RowMajor` and vice versa; `MatrixStride` comes from the buffer layout. One-row or one-column matrices and one-component vectors have no SPIR-V matrix type: they are vectors and scalars, and their indices disappear. - **Buffer layout** types (offset-decorated structs, strided arrays) are separate SPIR-V types from the plain types used in registers. Loading or storing a whole struct or array to a buffer converts member by member. - **Push constants.** A `[[vk::push_constant]]` block is a `Block` struct in the `PushConstant` storage class with no descriptor decorations. Its offsets follow the vector-relaxed std430 rules of a structured-buffer element (what DXC emits, checked against it by `PushConstantTests`), not the constant-buffer rules. - **Resources.** A constant buffer is a `Block` struct in `Uniform`; a structured or byte-address buffer is a `Block` struct holding one runtime array in `StorageBuffer` (`NonWritable` when read-only). Access chains into a storage buffer get the leading `0` for that struct. Textures and samplers are separate variables combined with `OpSampledImage` at each use. Storage images use the format from `[[vk::image_format]]` when the shader gives one (no read/write-without-format capability then; formats outside the core set add `StorageImageExtendedFormat`) and `Unknown` otherwise. Atomics on a texel go through `OpImageTexelPointer`; atomic scopes and barrier semantics match what DXC emits. A resource array is one variable of `OpTypeArray` (set and binding as for a single resource); a run-time index adds the matching `*ArrayDynamicIndexing` capability, and `NonUniformResourceIndex` adds `ShaderNonUniform`, the `*ArrayNonUniformIndexing` capability and `NonUniform` on the index, access chain, loaded handle and sampled image. An append or consume buffer's counter is a second storage buffer variable (`counter.var.`) in the buffer's set. - **Interface.** Every global variable the entry point references is listed on `OpEntryPoint`, as SPIR-V 1.4+ requires, inputs first, then outputs, then the rest in module order. Integer and double fragment inputs (and the vertex outputs that feed them) are `Flat`. Pixel shaders get `OriginUpperLeft`, plus `DepthReplacing` when they write depth. - **Stage variables.** Arrays and matrices (at least 2x2) become one variable per element or row, each with its own location, which occupies the same locations as DXC's single array-typed variable. Clip and cull distances are the exception: every `SV_ClipDistance`/`SV_CullDistance` variable of an entry point is a slice of one `ClipDistance` / `CullDistance` built-in array, ordered by semantic index (an array's elements in order), as DXC lays them out. - **Booleans in memory.** `OpTypeBool` is not allowed in uniform or storage memory, so `bool` members of buffers are 32-bit unsigned integers (0 or 1), converted to bool on load and back on store. A structured buffer's element stride is its size rounded up to its alignment (16 bytes for a `float3`). - **Control flow.** Blocks are emitted in reverse post-order of a walk that visits a header's merge and continue blocks first, which puts every merge after the blocks inside its construct. Unreachable blocks are dropped; a block without a terminator becomes `OpUnreachable`. - **Discard** is `OpKill`, valid in SPIR-V 1.5; `OpTerminateInvocation` needs a newer environment. - **Shifts** mask the count to the operand width, as HLSL does. `Ne` on floats is unordered (true for NaN). `precise` becomes `NoContraction`. ## Wave and quad intrinsics `GroupNonUniform*` instructions with Subgroup scope, chosen as DXC's output chooses them: `WaveReadLaneAt` is `OpGroupNonUniformShuffle`, `WaveActiveCountBits`/`WavePrefixCountBits` are `OpGroupNonUniformBallotBitCount` (Reduce / ExclusiveScan) of a ballot, the prefix functions are exclusive scans, `QuadReadAcross*` is `OpGroupNonUniformQuadSwap` and `QuadReadLaneAt` is `OpGroupNonUniformQuadBroadcast`. The lane index and count are the `SubgroupLocalInvocationId` and `SubgroupSize` input variables. The capabilities (`GroupNonUniform`, `Vote`, `Arithmetic`, `Ballot`, `Shuffle`, `Quad`) are added only for what the entry point uses; reflection lists the same classes (`EntryPointInfo.SubgroupOperations`, values equal to Vulkan's `VkSubgroupFeatureFlagBits`) so a host can check `subgroupSupportedOperations` before creating the pipeline. `WaveMatch` and `WaveMultiPrefix*` are front-end errors (OSH0350): partitioned operations exist in Vulkan only as a vendor extension. ## Specialization constants `[[vk::constant_id(N)]] const bool|int|uint|float Name = default;` is `OpSpecConstantTrue`/`OpSpecConstantFalse` (bool) or `OpSpecConstant` (the default's bits) with `OpDecorate SpecId N` and an `OpName`, written where it is first used, so an entry point carries only the constants it reaches. The ids are DXC's (a test compares the decorations with DXC's output for the corpus shaders). Expressions on constants are ordinary instructions; DXC emits `OpSpecConstantOp` for the initializer of a `static const` that depends on one, and the result is the same for every pipeline. The optimizer treats a constant as an unknown value, so branches on it stay. Pass the values in `VkSpecializationInfo` (`SpecializationValues.ToVulkan()`): every constant is a 4-byte value, a bool is a `VkBool32`. ## 16-bit and 64-bit types With `Enable16BitTypes`, `half`/`min16float`/`min10float` are `OpTypeFloat 16` and `min16int`/`min16uint` and the `int16_t` family are 16-bit integers (capabilities `Float16`, `Int16`, and for buffer members `StorageBuffer16BitAccess`, `UniformAndStorageBuffer16BitAccess`, `StoragePushConstant16`, `StorageInputOutput16`). `double` is `Float64`, `int64_t`/`uint64_t` are `Int64`, and 64-bit `Interlocked*` adds `Int64Atomics`. Buffer offsets, array strides and matrix strides equal DXC's (checked with `spirv-dis` on the corpus `types_*` cases). Without the switch, min-precision types are 32-bit and carry `RelaxedPrecision` on members and results; a `min16int` constant is cut to 16 bits while the arithmetic stays 32-bit (observed in DXC's output). Observed DXC choices followed: double arguments to float-only intrinsics are converted to float (we warn, OSH0361); unsuffixed float literals stay double until they meet a concrete type. One DXC behaviour is not copied: `(float)int64_t` goes through a 32-bit `SConvert` in DXC and truncates; we convert directly. `EntryPointInfo.RequiredFeatures` lists the device features the capabilities need. Note: on the RTX 4070 (driver 595) a 16-bit signed remainder with a negative dividend returns a different value from RADV, llvmpipe and Metal; the corpus case `types_int16` agrees between DXC and this compiler on every device because both emit `OpSRem`. ## Not supported yet (reported as OSH5002) Unbounded resource arrays (rejected by the front end), comparison or bit-casting of whole matrices, bool matrices, products involving one-row/column matrices. # The MSL backend `MslBackend.Emit(IrModule, entryPointName, MslOptions?, DiagnosticBag)` writes one entry point, and everything it reaches, as a single Metal Shading Language source file (language version 2.4 or newer; macOS 12 and iOS 15). `Compile` with `CompileTarget.Msl` returns one UTF-8 text blob per entry point. Code is in `src/OneShader/Msl/`. The same module and options always give the same text. Unsupported constructs are `OSH51xx` errors, never wrong code. ## Mapping decisions - **Functions.** `vertex`, `fragment` and `kernel` functions. The entry function's name is the HLSL name, except reserved words (`main` becomes `main0`, `fragment` becomes `fragment0`): `MslBackend.EntryFunctionName`. - **Push constants** are a `constant T& name [[buffer(n)]]` argument, where `n` is the next buffer index in declaration order like any constant buffer (recorded in reflection's `Msl` slot of the push block). The host fills it with `setVertexBytes`, `setFragmentBytes` or `setBytes`; the struct has the same byte layout as the SPIR-V block. - **Module-scope variables** do not exist in MSL. Stage variables, `static` and `groupshared` variables become locals of the entry function; every helper that touches one (directly or through its callees) receives it as an extra reference parameter. Resources are entry-function arguments and are passed to helpers the same way. - **Matrices.** HLSL `floatRxC` is MSL `floatRxC` (R columns of C components), so `m[i]` is row `i`, and `mul(a, b)` is `b * a`. Only float and half matrices exist in MSL; others are an error. - **Buffers keep the byte layout of the SPIR-V target.** Every buffer struct is emitted with `packed_*` vectors, padded array elements and explicit `uchar padN[]` members so each member sits at its reflected offset; `bool` is a 32-bit `uint`. Matrices are stored as lines (columns when column-major) and converted when loaded or stored. Compile the output with `-DONESHADER_CHECK_LAYOUT` to turn on `static_assert`s of every struct size and member offset against the reflected layout. - **Resource indices.** Buffers (constant, structured, byte-address), textures (including `Buffer`) and samplers each count up from 0 in declaration order over the whole file, and an array of N resources takes N indices. All entry points of a file see the same indices; they are the `Msl` `TargetSlot` in reflection. Direct indices are used rather than argument buffers, because they need no encoder objects on the host and are trivially deterministic. Limits: 31 buffers, 128 textures, 16 samplers per function (`OSH5105`). A vertex function with attributes reads them from buffer index `MslBackend.VertexBufferIndex` (30), one interleaved buffer using the reflected offsets and stride; resources may not use that index. - **Stage I/O.** Vertex inputs are a `[[stage_in]]` struct with `[[attribute(n)]]`; vertex outputs and fragment inputs use `[[user(locnN)]]` with `[[flat]]`, `[[center_no_perspective]]`, `[[centroid_perspective]]`, `[[sample_perspective]]`; fragment outputs use `[[color(n)]]`, `[[depth(any)]]`, `[[sample_mask]]`. `SV_VertexID`/`SV_InstanceID` are `[[vertex_id]]` and `[[instance_id]]`, which on Metal already include the draw's base vertex and base instance, as `VertexIndex`/`InstanceIndex` do in the SPIR-V target (`tools/metal-runner/probe-draw-ids.swift` shows it). Clip distances are `[[clip_distance]]`. Clip-space Y points up in Metal and down in Vulkan: flip with the viewport, the shader does not. - **Textures.** Textures used with comparison sampling become `depth*` types. Read-write textures get `read`, `write` or `read_write` access from their use. Array layers are split from the coordinate. Comparison samplers need their compare function set on the host sampler. - **Control flow.** Rebuilt from the IR's merge/continue markers; short-circuit branches without markers use post-dominators. The loop's continue block is repeated before each `continue`. - **Atomics** cast the address to `atomic_int`/`atomic_uint` with relaxed ordering. Compare-exchange loops over spurious failure. - **Resource arrays** (`Texture2D T[4]`, `SamplerState S[2]`, arrays of structured, byte-address and constant buffers) are `array, 4>`, `array`, `array` and `array` arguments; element `i` has Metal index `base + i`, which is what reflection reports as the slot of the array (its first index). `NonUniformResourceIndex(i)` is just `i`: Metal needs no hint. - **Append and consume buffers** use the hidden counter buffer as a second buffer argument, like the SPIR-V target: `Append` adds one with an atomic and stores at the old value, `Consume` subtracts one and loads. `GetDimensions` on them is the same error as on any structured buffer. ## Wave and quad intrinsics SIMD-group and quad-group functions (Metal 2.1 and later): `simd_sum`, `simd_product`, `simd_min`/`max`, `simd_and`/`or`/`xor`, `simd_prefix_exclusive_sum`/`product`, `simd_any`/`all`, `simd_ballot` (the 64-bit vote is split into the `uint4` HLSL returns), `simd_is_first`, `simd_shuffle` (`WaveReadLaneAt`), `simd_broadcast_first`, `quad_shuffle_xor` (`QuadReadAcross*`) and `quad_shuffle` (`QuadReadLaneAt`). `WaveGetLaneIndex` and `WaveGetLaneCount` read the `[[thread_index_in_simdgroup]]` and `[[threads_per_simdgroup]]` entry arguments once; helper functions get them as lifted parameters. `WaveActiveAllEqual` compares with `simd_broadcast_first`; the count-bits functions sum or scan `uint(condition)`. ## Specialization constants A constant is declared at file scope as `constant T Name [[function_constant(N)]];` (N is the HLSL `constant_id`, MSL 2.4 section 5.8.1) and read through `constant T Name_value = is_function_constant_defined(Name) ? Name : default;`, so a host that sets nothing gets the HLSL default. A function that declares function constants cannot build a pipeline state as it is, even when every default is wanted: create it with `makeFunction(name:constantValues:)`, using an empty `MTLFunctionConstantValues` for the defaults. `SpecializationValues.ToMetal()` gives index, `MTLDataType` number (float 3, int 29, uint 33, bool 53) and the 32-bit pattern for `setConstantValue(_:type:index:)`; a Metal bool is one byte. Only the constants an entry point reads are declared. ## 16-bit and 64-bit types `half`, `short`, `ushort`, `long` and `ulong` are the native MSL types (with `Enable16BitTypes`; without it `half` and min-precision are 32-bit). The only 64-bit atomics are `atomic_max_explicit` and `atomic_min_explicit` on an `atomic_ulong` in device memory, returning void (MSL spec 6.16.4.6, Apple silicon); the backend emits exactly that for unsigned `InterlockedMax`/`InterlockedMin` on a buffer whose result is unused. Any other 64-bit atomic, any use of `double`, and 64-bit integer vectors in buffers (no packed long vectors) are errors (OSH5102) at the line of the first use. Min-precision `RelaxedPrecision` has no MSL equivalent: values are computed in 32 bits. ## Not supported (errors) `double` (Metal has no 64-bit float); `SV_CullDistance`; `GetDimensions` on structured buffers (`ArrayLength`: Metal cannot query it); atomics on textures (Metal 3.1); `WaveMatch` and `WaveMultiPrefix*` (front-end error); the lane index and count in vertex shaders (Metal has no SIMD-group attribute for them); integer/bool matrices; 64-bit integer vectors in buffers; unbounded resource arrays; jumps out of several nested loops at once; comparison sampling on 1D/3D textures; texel offsets on cube/1D textures; overlapping `packoffset` members. Known deviations: `discard` in a helper that was not inlined calls `discard_fragment()` and returns, so the caller keeps running with writes suppressed; memory-only barriers are emitted as full threadgroup barriers (MSL 2.4 has no fence); signed integer overflow is undefined in MSL where HLSL wraps; Metal enables fast-math by default, so NaN and infinity handling is not guaranteed unless the host disables fast math. `precise` is honoured whatever the host's math mode. A function with precise arithmetic starts with `#pragma clang fp reassociate(off)` and `#pragma clang fp contract(off)` (no fused multiply-add, no reordered sums; measured with `xcrun metal`: fast math sets `fast` on every operation, these two pragmas remove `reassoc` and `contract`), and a precise division is `precise::divide`, because there is no pragma that stops `a / b` becoming `a * (1 / b)`. The flags `nnan`, `ninf`, `nsz`, `arcp` and `afn` stay set under fast math. `tools/run-metal-parity.sh run` with `ONESHADER_METAL_MATH=fast` and `safe` both match DXC on `lang_precise_math`; without the pragmas the fast run differs in 11 bytes. A `switch` whose cases fall through is written as a C++ `switch` with the same fall-through; values of the case's phis that arrive from the switch are assigned before it, and a `break` or `continue` inside a nested `if` in a case is written out. ## Checking `tools/validate-msl.sh ` compiles every `.metal` file in a directory with `xcrun -sdk macosx metal -c` (layout asserts on) and packs it with `metallib`. Produce the files with `ONESHADER_MSL_OUT= dotnet test tests/OneShader.Tests --filter-method '*MslBackendTests*'` (optimized files in ``, unoptimized in `/O0`). ## Running on a Metal GPU `tools/run-metal-parity.sh` runs every corpus case as MSL on a Metal device and compares it with DXC's SPIR-V on Vulkan. Three steps share one exchange directory: `export` (a machine with Vulkan and DXC) writes the MSL, the inputs and the reference outputs; `run` (a Mac with Xcode, no .NET needed) builds `tools/metal-runner` and executes each case with Metal API and shader validation on, so a validation message fails the case; `compare` prints the per-case table. Allowed differences are in `corpus/metal-tolerances.txt`; everything else must match as the Vulkan parity suite requires. The runner draws with a viewport of negative height. Clip-space Y points up in Metal and down in Vulkan, and the corpus shaders are written for Vulkan; the flipped viewport puts the same vertex positions on the same pixels (including `SV_Position`). The mirrored viewport reverses the winding Metal sees, so Vulkan's counter-clockwise front face is also counter-clockwise on Metal (back-face culling on `pbr_forward` shows it). Depth range and the texture origin already agree. Vertex buffer binding N is Metal buffer 30 - N, described to Metal with a vertex descriptor. # The WGSL backend `WgslBackend.Emit(IrModule, entryPointName, WgslOptions?, DiagnosticBag)` writes one entry point, and everything it reaches, as WGSL text (`CompileTarget.Wgsl` gives one UTF-8 blob per entry point). It returns `null` and reports `OSH52xx` errors for constructs WGSL cannot express; it never throws for unsupported input. Output is deterministic. Code is in `src/OneShader/Wgsl/`; naga validates every corpus entry point (`tests/OneShader.Tests/WgslBackendTests.cs`). ## Mapping decisions - **Push constants** (WebGPU core has none) are a `var` of 16-byte rows like a constant buffer, with the SPIR-V block's byte layout. It is bound at the highest bind group the shader uses and the binding after that group's highest one, so it cannot collide with the shader's own bindings; the slot is in reflection and in the generated `Bindings.PushConstants`. The host binds a uniform buffer of at least the block size rounded up to 16 bytes. - **Buffers are raw words, so layouts equal the SPIR-V target's.** A constant buffer is `array, N>` (WGSL 13.4.1: uniform arrays have a 16-byte stride, one vec4 per row); structured and byte-address buffers are `array` (`array>` when the entry point uses atomics on them; a buffer the entry point never writes is `read`). Every access computes its byte offset from the reflected layout and bitcasts; vectors inside one row are read as a swizzle. Signed atomic min/max on a buffer and compare-exchange use small helper functions. `WgslLayoutTests` checks the word indices of every scalar, vector and matrix in every corpus buffer against the reflected offsets. - **Matrices are transposed**, as in SPIR-V: HLSL `floatRxC` is `matRxC` (R columns of C components), `m[i]` is row i, and every `mul` swaps its operands. Matrix products with one-row or one-column matrices are not supported. - **Bindings:** `@group(set) @binding(binding)`, equal to the Vulkan set and binding. - **Stage I/O:** inputs and outputs are module-level private variables copied from the entry parameters and into the return struct. Arrays and matrices are split one location per element or row. Integer fragment inputs and vertex outputs are `flat`. - **Control flow** is rebuilt from the IR's structure markers: `if`, `switch`, `loop` with `continuing` (a do-while test becomes `break if`). A `break` out of a loop from inside a `switch` goes through a flag variable. Phis are variables assigned on each edge. WGSL has no fall-through, so a `switch` case that runs into the next one gets that case's statements again after its own (a chain of n cases that fall through each other is written out in full, so the text grows with n squared); a `default` in the middle becomes the last clause. - **Entry names:** the HLSL name, with anything that is not a letter, digit or underscore turned into an underscore and a trailing underscore added to names WGSL reserves or predeclares (`loop` becomes `loop_`): `WgslBackend.EntryFunctionName`, which reflection and the generated `WgslEntryPoint` use. - **Textures:** a texture used with a comparison sampler is a depth texture; storage textures need `[[vk::image_format]]` and take `read`, `write` or `read_write` from how the entry point uses them. - **Derivatives:** pixel shaders that sample or take derivatives start with `diagnostic(off, derivative_uniformity);`, since HLSL allows them under non-uniform control flow. ## Wave and quad intrinsics WebGPU's optional `subgroups` feature. A module that uses any wave or quad intrinsic starts with `enable subgroups;` and `diagnostic(off, subgroup_uniformity);` (HLSL does not require wave operations to run in uniform control flow), and the lane index and count are the `subgroup_invocation_id` and `subgroup_size` built-ins. The host must create the device with the `subgroups` feature: reflection (`ShaderReflection.UsesSubgroups`, `EntryPointInfo.SubgroupOperations`) and `PipelineLayout.TryToWebGpu` (`RequiredFeatures`) say so. Mapping: `subgroupAdd`/`Mul`/`Min`/`Max`/`And`/`Or`/`Xor`, `subgroupExclusiveAdd`/`Mul` for the prefix functions, `subgroupBallot`, `subgroupAny`/`All`, `subgroupElect`, `subgroupShuffle`, `subgroupBroadcastFirst`, `quadSwapX`/`Y`/`Diagonal`, `quadBroadcast`; `WaveActiveAllEqual` compares with `subgroupBroadcastFirst` and the count-bits functions add `u32(condition)`. naga's WGSL reader does not parse the subgroups extension yet, so these modules are validated by the browser's compiler (Tint) in the WebGPU parity run instead. ## Specialization constants A constant is `@id(N) override Name: T = default;` (WGSL section 7.3; N is the HLSL `constant_id`, which must be at most 65535 here, otherwise OSH5202). Only the constants an entry point reads are declared. The pipeline's `constants` record is keyed by the decimal id, which is also the key when an override has a name: `SpecializationValues.ToWebGpu()` builds it, with a bool as 0 or 1. WebGPU rejects a key the shader does not declare, so use the constants of the entry point you are building (`EntryLayout.SpecConstants`, or the generated `SpecConstantIds`). The `naga` command-line tool writes overrides into SPIR-V as plain constants at their defaults, so overrides are verified in the browser run rather than through naga. ## 16-bit floats With `Enable16BitTypes`, `half` and `min16float` are `f16` and the module starts with `enable f16;` (WebGPU's optional `shader-f16` feature: reflection `DeviceFeatures.Float16` and `PipelineLayout.TryToWebGpu` `RequiredFeatures` say when to request it). Buffers stay raw 32-bit words: a half is read as one half of `unpack2x16float(word)` and written with `pack2x16float`; because the other half of a word may belong to a neighbour written by another invocation, a buffer written with halves is declared `array>` and a store clears and sets only its own 16 bits with `atomicAnd`/`atomicOr`. WebGPU does not require half arithmetic to round exactly as Vulkan does, so results can differ in the last bit of a half (the parity case allows 0.4%). There are no 16-bit integers. ## Not available (diagnostic OSH5202) clip and cull distances (optional WebGPU feature); `CalculateLevelOfDetail`; wave and quad intrinsics in vertex shaders (WebGPU's subgroup functions and built-ins exist only in compute and fragment shaders); resource arrays (WebGPU has no binding arrays: one error per array, which suggests separate resources or a `Texture2DArray`; `NonUniformResourceIndex` is just its index); `Buffer`/`RWBuffer`; append and consume buffers; atomics on textures; `double`, 64-bit integers and 16-bit integers (each reported at the line of its first use); integer and bool matrices; `PrimitiveID`, layer and viewport index; recursion; storage textures without a format or read-written in a format other than r32float/r32uint/r32sint; bool stage variables; texel offsets on cube textures; control flow that is not structured. ## Known differences WGSL lets an implementation fuse a multiply and an add, so float results can differ in the last bits from the SPIR-V target's; in the parity run against DXC the `particles` case differs in the last bits on one driver version and matches on others. `precise` has no WGSL equivalent: the arithmetic compiles, with warning OSH5205 once per entry point, and the result may differ from the other targets in the last bits. Memory-only barriers become full barriers and must sit in uniform control flow. `isnan`, `isinf` and `isfinite` test exponent bits, which WGSL implementations may assume never fire. ## Checking Every corpus file, optimized and not, either validates with `naga` or reports one of the errors above (the table in `WgslBackendTests`). `ONESHADER_PARITY_TARGET=wgsl tools/run-parity.sh` runs the GPU comparison against DXC on the WGSL output translated to SPIR-V by `naga`; cases that need an unavailable construct are skipped and listed as `unsupported-on-target` in the parity table. # The intermediate representation One typed IR sits between the front end and every backend (SPIR-V, MSL, WGSL). Its types live in `src/OneShader/IR/`; the HLSL type system it shares is in `src/OneShader/Types/`. ## Shape: SSA values, structured control flow, memory for locals - **Values are SSA.** Every instruction that yields a value defines it once; operands are `IrValue`s (`Constant`, `Argument`, `GlobalVariable`, `Instruction`). `Phi` exists, but the lowering from HLSL starts from the simplest form: each HLSL local is an `Alloca` in the entry block accessed with `Load`/`Store`, and a later promotion pass can turn those into SSA with phis. Backends must accept both forms. - **Control flow is a CFG that keeps its structure.** Blocks end in exactly one terminator (`Branch`, `CondBranch`, `Switch`, `Return`, `Kill`, `Unreachable`). Header blocks carry the same markers SPIR-V needs: - `StructureKind.Selection`: ends in `CondBranch` or `Switch`; `MergeBlock` is where the arms rejoin. - `StructureKind.Loop`: ends in `Branch`; `MergeBlock` is the exit, `ContinueBlock` is the block that branches back to the header. `continue` is a branch to the continue block, `break` a branch to the merge block of the innermost loop or switch. - `ControlHint` carries `[unroll]`, `[loop]`, `[flatten]`, `[branch]`. Passes must keep these markers valid (the verifier checks they exist and point into the function). Because they are explicit, SPIR-V gets its `OpSelectionMerge`/`OpLoopMerge` directly and the text backends can rebuild `if`/`for`/`switch` without a general structurizer. - **Addressing is logical.** `PointerType(AddressSpace, pointee)` values come only from `Alloca`, globals and `AccessChain`. There is no pointer arithmetic or pointer-to-pointer. `AccessChain` takes a base pointer plus indices (constant `int` for struct members, any integer for arrays and vector components). - **Types are `HlslType`.** Scalars, vectors, matrices, arrays, structs, resources and pointers are the same classes the front end uses, so lowering does not translate types. Matrix constants are composites of rows. ## Module layout `IrModule` holds `Globals`, `Functions` and `EntryPoints` in source order (the dump and every backend follow that order). An `EntryPoint` names a no-parameter void function and lists its `Inputs` and `Outputs`: stage variables are module globals (`GlobalKind.Input`/`Output`) with a `Location` or `BuiltIn`; the lowering wraps the HLSL entry function that takes and returns structs. Resources are globals of pointer-to-`ResourceType` in `AddressSpace.UniformConstant` (textures, samplers, typed buffers), `Uniform` (constant buffers) or `StorageBuffer` (structured and byte-address buffers), with `Set`/`Binding` and, for buffers, the `TypeLayout` from `BufferLayout`. A `cbuffer` is a `ConstantBuffer` resource whose element is a struct of its members. Access chains into `ConstantBuffer` start inside `T` (the first index selects a member of `T`); chains into `StructuredBuffer` start with the element index, then continue inside `T`; a chain with one index into a byte-address buffer addresses a 32-bit word (used for atomics). A `SpecConstant` is a scalar value like a `Constant` (it can be an operand anywhere a scalar is) with an id and a default, listed in `IrModule.SpecConstants`. It is deliberately not a `Constant`: no pass folds it, so code that depends on it survives the optimizer. The dump prints `spec ID NAME = default` lines and `spec(ID NAME)` operands. `CompileOptions.SpecializationOverrides` replaces chosen ones by real constants after lowering. ## Opcode conventions Arithmetic and comparisons are component-wise and typed by their operands; signedness and float-versus-integer come from the operand type, there are no separate `fadd`/`iadd` opcodes. `Mul` on matrices is component-wise; the matrix product is `MatMul`. `Precise` on an instruction forbids reassociation, fusing and algebraic rewrites. Resource operands of texture and buffer opcodes are pointers (the global, or an access chain into a resource array), never loaded handles. Texture operands, in order: | Opcode | Operands | Immediates | |---|---|---| | `TextureSample` | texture, sampler, coordinate | constant texel offset (0 to 3 values) | | `TextureSampleBias` / `Level` | texture, sampler, coordinate, bias or level | offset | | `TextureSampleGrad` | texture, sampler, coordinate, ddx, ddy | offset | | `TextureSampleCompare` / `CompareLevelZero` | texture, comparison sampler, coordinate, reference | offset | | `TextureGather` | texture, sampler, coordinate | component, then offset | | `TextureGatherCompare` | texture, comparison sampler, coordinate, reference | offset | | `TextureLoad` | texture, integer coordinate, optional level (or sample index when multisampled) | offset | | `TextureStore` | texture, integer coordinate, value | | | `TextureQuerySize` | texture, optional level | | | `TextureQueryLevels`, `TextureQuerySamples` | texture | | | `TextureQueryLod` | texture, sampler, coordinate | | | `ByteAddressLoad` / `Store` | buffer, byte offset (, value) | | `AtomicRmw` and `AtomicCompareExchange` take a pointer to the target (a groupshared variable, a buffer element, or a one-index access chain into a byte-address buffer, which addresses a 32-bit word). On a read-write texture the first operand is the texture pointer and the second the integer coordinate; the values follow (`AtomicRmw`: value, `AtomicCompareExchange`: comparand, then new value). Both return the original value. `Barrier`'s immediate is a `BarrierFlags` combination. A global's `ImageFormat` (a name from `ImageFormats`, the `[[vk::image_format]]` value) says how a storage texture is laid out; without it the format is unspecified. Array textures put the layer in the last coordinate component. `CallIntrinsic` takes the operands of the HLSL intrinsic in HLSL order; intrinsics with `out` parameters (`modf`, `frexp`, `sincos`) take pointers. Functions with resource-pointer parameters must be inlined before a backend runs. ## Dump format `IrPrinter.Print` is the textual form tests compare. Values are numbered in print order per function (`%0`, `%1`; parameters are `%arg0`), blocks by position (`bb0`), so output does not depend on object identity, hash order or names. Source names appear only in `;` comments. Example: ``` func @f(%arg0: int) -> int { bb0: ; (entry) %0 = alloca : ptr ; sum store %0, 0:int br bb1 bb1: ; loop merge=bb5 continue=bb4 hint=dontunroll (for.header) ... ``` ## Verifier `IrVerifier.Verify` returns messages and never throws. It checks block shape (one terminator, last), parent links, operand ownership and in-block ordering, phi placement, the types of loads, stores, arithmetic, comparisons, selects, branches, calls and phis, and the structure markers. For reachable blocks it also checks that every definition dominates its uses and that each phi names exactly its block's predecessors (`Cfg` in `src/OneShader/IR/Cfg.cs` provides predecessors, reverse post-order, dominators and dominance frontiers). Tests run it after building and after every optimizer pass. ## Optimizer `CompileOptions.Optimize` (default true) runs `Optimizer.Run` (`src/OneShader/Optimization/`) between lowering and the backends. Reflection is built from the unoptimized module, so it does not change with the optimizer. 1. **Inlining** of every call, callees first. A callee whose only `ret` is outside every structured construct is spliced in directly. Any other callee is wrapped in a loop that runs once: a `ret` stores the result and breaks; a `ret` inside an inner loop or switch sets a flag and breaks, and that construct's merge block tests the flag and breaks outward, so each branch is a legal SPIR-V break. Recursive calls are left alone, and growth is capped. 2. Per function, repeated until nothing changes: **SROA** (structs and small arrays with constant indices become one variable per element), **mem2reg** (scalars, vectors and matrices become SSA values; phis are placed on the dominance frontier, pruned by liveness), **combine** (constant folding and exact algebraic identities; `precise` instructions are only folded), **CSE** along the dominator tree (pure instructions, plus loads from constant buffers, resources and stage inputs), **DCE** (mark and sweep, plus stores to variables nobody reads) and **CFG simplification** (constant branches, unreachable blocks, straight-line merging). 3. Folding is limited to results that are the same on every implementation: wrapping integer arithmetic, masked shifts, float add/sub/mul/div/rem in the operand's own precision. Division by zero, `INT_MIN / -1`, NaN and infinite or subnormal operands and results, out-of-range float-to-int conversions and transcendental functions are left to the GPU. 4. Merge and continue blocks that a live header names are kept even when unreachable (reduced to `unreachable` or a branch to the header); headers, merge blocks and continue blocks are never merged away. If a pass fails, `Optimizer.Run` returns a message instead of throwing and `Compile` falls back to the unoptimized module with warning OSH6001. Setting `Optimizer.VerifyEachPass` (tests do) verifies the IR after every pass. ## Decisions worth knowing - One `Instruction` class for all opcodes, so passes can rewrite in place; `IrBuilder` provides typed helpers. There are no use lists: passes build a user map per run (`IrUtil.BuildUsers`) or collect replacements in a map and apply them in one sweep (`IrUtil.ApplyReplacements`). - Constants keep exact bits (`Constant.Bits`), so folding wraps integers as HLSL does and keeps float results in the operand's own precision. - Block names and value names are hints and never take part in equality or output ordering. ## Lowering conventions (front end) - An HLSL entry point becomes two functions: `Name_body`, the user function with its original parameters, and a wrapper named `Name` that loads the input globals into a struct, calls the body and stores the outputs. Only the wrapper is an `EntryPoint`. - Resource arrays are globals whose value type is an `ArrayType` of the resource (`ResourceCount` set); an element is an `AccessChain` with the index. Constant-buffer members are `AccessChain(global, type, memberIndex)`. A function that takes a resource is compiled once per distinct resource it is called with; a run-time index into a resource array is passed to that copy as an ordinary `int` argument. - `NonUniformResourceIndex(i)` is the intrinsic `NonUniformIndex`: it yields `i` and tells the SPIR-V backend to decorate the access chain, the loaded handle and the sampled image built from it as `NonUniform`. - An append or consume buffer comes with a second global, `counter.var.`, a `RWByteAddressBuffer` in the same set. `Append` is `AtomicRmw(Add, counter[0], 1)` followed by a store at the returned slot; `Consume` is an `AtomicRmw(Add, counter[0], 0xFFFFFFFF)` followed by a load at the returned value minus one. - Non-constant `static` initializers run in a function `init_globals` that each wrapper calls first. - `&&`, `||` and non-pure `?:` are control flow with a `Phi`; pure `?:` is `Select`. - Initializer lists follow HLSL's flat rules: a run of values fills an aggregate in order (a vector-valued expression supplies several scalars) and a braced sub-list starts a new element. - Dead code after `return`, `break`, `continue` and `discard` is not lowered; merge blocks nothing reaches end in `Unreachable`. # WebGPU execution parity `tools/run-webgpu-parity.sh` runs the WGSL output of every corpus case in headless Chromium on a real WebGPU adapter and compares the results with DXC's SPIR-V running on Vulkan. ``` ONESHADER_TOOLS= ONESHADER_ARTIFACTS= tools/run-webgpu-parity.sh ``` What it does: 1. `tests/OneShader.ParityTests/WebGpuParityTests.cs` compiles each case's entry points with `CompileTarget.Wgsl`, runs the DXC-compiled SPIR-V on Vulkan (validation layer on) as the reference, and writes one JSON job per case (WGSL, buffer and texture bytes, pipeline description). 2. `tools/webgpu/run.py` starts Chromium through Playwright (a `uv` environment under `ONESHADER_WEBGPU_HOME`, default `~/.cache/oneshader-webgpu`; Playwright downloads its own Chromium) and serves `tools/webgpu/index.html`. `runner.js` creates the buffers, textures and samplers, builds bind group layouts from the WGSL's own declarations, runs the compute dispatches or the offscreen draw, and reads the outputs back. Every job runs inside validation, out-of-memory and internal error scopes plus an uncaptured-error listener, and shader compilation messages are checked; any of these fails the case. 3. The test compares outputs like the other parity runs: integer and other buffers byte for byte, 8-bit images within the case's `tolerance`, float images within `ftolerance`. Output buffers marked `float=X` in a case file (32-bit floats whose math WebGPU only specifies loosely, such as logarithms) compare within `|a-b| <= X * (1 + |a|)`. Cases the backend deliberately rejects (diagnostic OSH52xx) are skipped with the diagnostic. The browser must get a hardware adapter: the driver exits with an error when Chromium offers only a software one. For debugging the tooling without a GPU, set `ONESHADER_ALLOW_SOFTWARE_ADAPTER=1` and `ONESHADER_BROWSER_ARGS_ONLY="--enable-unsafe-webgpu --no-sandbox"` (SwiftShader), and `ONESHADER_VK_DEVICE=llvmpipe` for the reference; rasterization and filtering then differ between the two software implementations, so image mismatches are expected. `tools/webgpu/probe.py` prints the adapter a launch gets. Chromium flags that give a hardware adapter on Linux: `--enable-unsafe-webgpu --enable-features=Vulkan --use-angle=vulkan --ignore-gpu-blocklist` (see `tools/webgpu/browser.py`). ## Orientation Clip-space y points up in WebGPU, like D3D and Metal. The Vulkan reference draws with a negative-height viewport and clockwise front faces (`VkRunner.Run(..., yUp: true)`), so both sides see the same clip coordinates and framebuffer rows (`SV_Position`). The corpus meshes wind counter-clockwise for Vulkan's y-down space; on a y-up target the same vertices need `frontFace: 'cw'` to keep `SV_IsFrontFace` and back-face culling identical, which is what the runner sets. A shader and mesh authored for y-up conventions would use `ccw` instead. A shader written for Vulkan's y-down clip space draws upside down on WebGPU unless its projection accounts for it; the backend does not flip. ## Options and limits The device is requested with default limits and no optional features, which is what a portable web build gets. `ONESHADER_WEBGPU_CASES=a,b` runs a subset. # Performance OneShader compiles a shader about 2.4 times faster than DXC when both run as native processes, and about 14 times slower when it runs as an ordinary .NET process that has to start the runtime and JIT-compile the compiler first. Both are measured below, on the same 68 entry points, with the compiler's SPIR-V compared against DXC's SPIR-V. This page prints all of it: the tables, the losses and the method. Every number comes from the files in [`bench/results/`](https://github.com/recursive-ai-oss/oneshader-csharp/tree/8896e63be3a4f246295427451f15e4fcfb5535cf/bench/results), measured on one workstation (see [Method](#method)). Treat them as local measurements, not promises for your hardware. ## Results at a glance One shader compile per process, for each of the 68 (file, entry point) pairs the repository's execution-parity cases compile (25 2D-renderer, 24 3D, 11 compute, 8 language-feature entry points). DXC is the released binary run with `-spirv`; its default optimization. OneShader is a small command-line wrapper around `ShaderCompiler.CompileFile`, built three ways: as a NativeAOT binary, as a framework-dependent ReadyToRun build, and as a plain framework-dependent (JIT) build. Time is wall-clock from creating the process to its exit, with the blob written to disk. The ratio is DXC's time divided by ours per entry point, summarized by the geometric mean; a time within 5% of DXC's is a tie. Bar chart: how many times faster or slower OneShader's compile is than DXC's, per task, log scale. NativeAOT is about 2.4 times faster warm and 2.7 times faster cold. ReadyToRun is about 5 to 6 times slower and the JIT build 13 to 14 times slower. The chart is drawn from the results file by [`bench/make-report.py`](https://github.com/recursive-ai-oss/oneshader-csharp/blob/8896e63be3a4f246295427451f15e4fcfb5535cf/bench/make-report.py). | Task (one process per entry point) | Entry points | Time vs DXC (geometric mean) | Fastest | Slowest | Entry points slower than DXC | |---|--:|---|---|---|--:| | NativeAOT, warm: 2D renderer | 25 | 2.46× faster | 2.82× faster | 2.09× faster | 0 | | NativeAOT, warm: 3D (PBR set) | 24 | 2.41× faster | 3.14× faster | 2.13× faster | 0 | | NativeAOT, warm: compute | 11 | 2.49× faster | 2.81× faster | 2.21× faster | 0 | | NativeAOT, warm: language features | 8 | 2.30× faster | 2.72× faster | 2.08× faster | 0 | | NativeAOT, warm: all | 68 | 2.43× faster | 3.14× faster | 2.08× faster | 0 | | NativeAOT, cold: all | 68 | 2.68× faster | 3.57× faster | 1.51× faster | 0 | | ReadyToRun, warm: all | 68 | 6.10× slower | 3.21× slower | 7.74× slower | 68 | | ReadyToRun, cold: all | 68 | 5.30× slower | 2.78× slower | 7.00× slower | 68 | | JIT, warm: all | 68 | 14.3× slower | 7.28× slower | 17.7× slower | 68 | | JIT, cold: all | 68 | 12.9× slower | 6.73× slower | 16.3× slower | 68 | "Warm" is the median of 15 repeated runs per entry point; "cold" is the first run of each entry point right after the corpus was written to a fresh directory. ### The losses - **Started as a .NET process, OneShader is slower than DXC on every entry point.** The JIT build is 14.3× slower warm (113 ms median per entry point against DXC's 7.6 ms) and 12.9× slower cold; ReadyToRun is 6.1× and 5.3× slower (48 ms median). Almost all of that is starting the runtime and compiling or loading the compiler's code, not compiling the shader: the same compile inside a running process takes 0.34 ms at the median ([in-process](#in-process-compile-time)). The NativeAOT build is the one that wins, so a command-line tool that compiles shaders should ship NativeAOT; the MSBuild task runs inside the build's own long-lived .NET process, where the runtime start-up is paid once per build, not once per shader (that case, the first compile in a fresh process, is not measured here). - **The SPIR-V is larger than DXC's with debug names on, which is the default.** 57 of 68 entry points are larger, 46,620 words against DXC's 44,126 (+5.7%); the worst are `lang/sizeof_static.hlsl:CSMain` (+90%), `3d/skinned.hlsl:VSMain` (+60%) and `3d/stage_io.hlsl:VSMain` (+36%). With names off (`CompileOptions.KeepNames = false`) the total is 40,258 words, 8.8% below DXC, with 61 entry points smaller and 7 larger. DXC emits no names unless asked, so the second comparison is the like-for-like one; the first is what the compiler gives you by default ([table](#output-size)). - **The optimizer is the largest cost of a compile.** With it off, SPIR-V takes 0.13 ms at the median instead of 0.34 ms (2.5× less); the worst file, `3d/bloom_downsample.hlsl`, takes 1.18 ms with the optimizer and 0.22 ms without (5.5×). That is the price of the smaller output; see the in-process table. - **MSL and WGSL are not on the chart.** DXC has no such outputs, so there is nothing to compare; their in-process times are in the table. They also refuse some corpus shaders, which are left out of those columns: MSL refuses `compute/append_consume.hlsl`; WGSL refuses that file, `3d/stage_io.hlsl` and `lang/resource_arrays.hlsl` (the diagnostics are in [`bench/results/inproc/`](https://github.com/recursive-ai-oss/oneshader-csharp/tree/8896e63be3a4f246295427451f15e4fcfb5535cf/bench/results/inproc)). - **No in-process comparison with DXC.** DXC is timed as a process only. Its in-process cost is not measured, so the in-process numbers below stand alone. ## In-process compile time `ShaderCompiler.Compile` called on the source text in a warm process, every entry point of a file in one call, SPIR-V with the default optimizer unless the column says otherwise. BenchmarkDotNet 0.15.8 with the memory diagnoser, in-process toolchain, 3 warm-up and 12 measured iterations of at least 150 ms each (10 to 12 kept after outlier removal); the mean in milliseconds, then the standard deviation in brackets, and the managed allocation per call in KiB. Includes (`3d/brdf.hlsli`) are read from disk during the compile; the shader text is read once beforehand. Over the 41 files: the median file compiles in 0.34 ms and allocates 0.85 MiB; the slowest is `2d/renderer/ui.hlsl` at 2.0 ms and 4.2 MiB. MSL and WGSL cost about 3 to 5% more than SPIR-V. The first row, `2d/blit.hlsl` with SPIR-V, was the first benchmark the process ran and shows it (a standard deviation of 34% of its mean; the same file with MSL, run later, takes 0.12 ms): read that one cell as a warm-up artifact. | File | SPIR-V ms | alloc KiB | SPIR-V, optimizer off ms | alloc KiB | MSL ms | alloc KiB | WGSL ms | alloc KiB | |---|--:|--:|--:|--:|--:|--:|--:|--:| | `2d/blit.hlsl` | 0.224 (±0.076) | 335 | 0.055 (±0.001) | 195 | 0.124 (±0.000) | 353 | 0.125 (±0.001) | 360 | | `2d/blur.hlsl` | 0.264 (±0.000) | 689 | 0.106 (±0.000) | 346 | 0.281 (±0.000) | 721 | 0.280 (±0.000) | 733 | | `2d/color_grade.hlsl` | 0.278 (±0.003) | 728 | 0.116 (±0.000) | 383 | 0.292 (±0.001) | 763 | 0.294 (±0.001) | 789 | | `2d/msdf_text.hlsl` | 0.313 (±0.000) | 822 | 0.130 (±0.000) | 429 | 0.335 (±0.001) | 873 | 0.334 (±0.001) | 893 | | `2d/renderer/quad.hlsl` | 0.654 (±0.001) | 1704 | 0.256 (±0.001) | 859 | 0.692 (±0.001) | 1789 | 0.695 (±0.003) | 1835 | | `2d/renderer/tile.hlsl` | 0.551 (±0.001) | 1437 | 0.227 (±0.003) | 744 | 0.586 (±0.001) | 1508 | 0.586 (±0.001) | 1543 | | `2d/renderer/ui.hlsl` | 2.018 (±0.007) | 4268 | 0.666 (±0.004) | 1684 | 2.131 (±0.020) | 4591 | 2.194 (±0.022) | 4668 | | `2d/sdf_text.hlsl` | 0.248 (±0.000) | 682 | 0.107 (±0.000) | 368 | 0.269 (±0.001) | 728 | 0.266 (±0.000) | 747 | | `2d/sprite.hlsl` | 0.163 (±0.000) | 446 | 0.077 (±0.001) | 248 | 0.179 (±0.000) | 478 | 0.178 (±0.000) | 502 | | `2d/tilemap.hlsl` | 0.269 (±0.000) | 697 | 0.116 (±0.001) | 374 | 0.288 (±0.001) | 738 | 0.288 (±0.000) | 760 | | `2d/ui_rect.hlsl` | 0.666 (±0.005) | 1707 | 0.242 (±0.001) | 785 | 0.707 (±0.001) | 1821 | 0.710 (±0.001) | 1909 | | `3d/bloom_downsample.hlsl` | 1.181 (±0.002) | 2941 | 0.215 (±0.001) | 741 | 1.187 (±0.003) | 3101 | 1.191 (±0.004) | 3135 | | `3d/bloom_upsample.hlsl` | 0.343 (±0.006) | 868 | 0.133 (±0.000) | 446 | 0.351 (±0.000) | 915 | 0.350 (±0.000) | 939 | | `3d/clustered_lights.hlsl` | 0.953 (±0.011) | 2332 | 0.355 (±0.002) | 1060 | 1.002 (±0.001) | 2467 | 0.997 (±0.004) | 2478 | | `3d/ibl.hlsl` | 0.649 (±0.001) | 1662 | 0.318 (±0.003) | 1014 | 0.660 (±0.001) | 1793 | 0.646 (±0.001) | 1760 | | `3d/pbr_forward.hlsl` | 1.403 (±0.029) | 3143 | 0.561 (±0.007) | 1394 | 1.423 (±0.005) | 3345 | 1.402 (±0.002) | 3361 | | `3d/shadow_cascades.hlsl` | 1.104 (±0.003) | 2679 | 0.363 (±0.002) | 1080 | 1.175 (±0.003) | 2833 | 1.162 (±0.002) | 2862 | | `3d/shadow_depth.hlsl` | 0.100 (±0.000) | 254 | 0.048 (±0.000) | 143 | 0.109 (±0.000) | 274 | 0.110 (±0.000) | 300 | | `3d/skinned.hlsl` | 0.334 (±0.001) | 888 | 0.133 (±0.001) | 454 | 0.345 (±0.001) | 947 | 0.359 (±0.001) | 1083 | | `3d/ssao.hlsl` | 0.528 (±0.009) | 1299 | 0.198 (±0.003) | 649 | 0.542 (±0.001) | 1370 | 0.540 (±0.001) | 1407 | | `3d/stage_io.hlsl` | 0.367 (±0.001) | 962 | 0.147 (±0.001) | 484 | 0.387 (±0.001) | 1006 | refused | – | | `3d/tonemap.hlsl` | 0.502 (±0.001) | 1179 | 0.189 (±0.001) | 556 | 0.522 (±0.001) | 1218 | 0.516 (±0.001) | 1228 | | `compute/append_consume.hlsl` | 0.162 (±0.002) | 417 | 0.073 (±0.001) | 228 | refused | – | refused | – | | `compute/cluster_build.hlsl` | 0.419 (±0.001) | 1107 | 0.129 (±0.001) | 412 | 0.440 (±0.001) | 1155 | 0.441 (±0.002) | 1205 | | `compute/cull_instances.hlsl` | 0.353 (±0.007) | 836 | 0.125 (±0.002) | 373 | 0.368 (±0.000) | 876 | 0.362 (±0.001) | 871 | | `compute/depth_bounds.hlsl` | 0.209 (±0.000) | 562 | 0.093 (±0.001) | 307 | 0.223 (±0.002) | 594 | 0.220 (±0.001) | 599 | | `compute/histogram.hlsl` | 0.152 (±0.000) | 409 | 0.072 (±0.001) | 228 | 0.161 (±0.001) | 426 | 0.156 (±0.000) | 423 | | `compute/ibl_prefilter.hlsl` | 0.655 (±0.006) | 1632 | 0.232 (±0.002) | 753 | 0.697 (±0.001) | 1723 | 0.687 (±0.001) | 1734 | | `compute/image_filter.hlsl` | 0.314 (±0.000) | 803 | 0.126 (±0.002) | 404 | 0.323 (±0.000) | 823 | 0.329 (±0.001) | 853 | | `compute/light_cull.hlsl` | 0.264 (±0.000) | 690 | 0.108 (±0.002) | 355 | 0.278 (±0.000) | 724 | 0.272 (±0.001) | 731 | | `compute/particles.hlsl` | 0.291 (±0.000) | 741 | 0.108 (±0.002) | 370 | 0.312 (±0.000) | 803 | 0.313 (±0.001) | 831 | | `compute/prefix_sum.hlsl` | 0.164 (±0.000) | 433 | 0.068 (±0.001) | 221 | 0.169 (±0.001) | 442 | 0.171 (±0.000) | 457 | | `compute/stage_layouts.hlsl` | 0.210 (±0.000) | 592 | 0.095 (±0.000) | 344 | 0.213 (±0.000) | 606 | 0.217 (±0.001) | 631 | | `lang/inheritance.hlsl` | 0.358 (±0.001) | 920 | 0.125 (±0.002) | 409 | 0.367 (±0.000) | 937 | 0.364 (±0.000) | 944 | | `lang/members.hlsl` | 0.397 (±0.003) | 1035 | 0.124 (±0.000) | 411 | 0.405 (±0.001) | 1050 | 0.405 (±0.003) | 1071 | | `lang/operators.hlsl` | 0.657 (±0.002) | 1874 | 0.224 (±0.003) | 778 | 0.659 (±0.002) | 1872 | 0.660 (±0.001) | 1882 | | `lang/resource_arrays.hlsl` | 0.269 (±0.002) | 697 | 0.123 (±0.001) | 399 | 0.274 (±0.001) | 711 | refused | – | | `lang/resource_locals.hlsl` | 0.278 (±0.001) | 759 | 0.105 (±0.002) | 357 | 0.285 (±0.000) | 768 | 0.287 (±0.001) | 794 | | `lang/sizeof_static.hlsl` | 0.227 (±0.000) | 646 | 0.083 (±0.001) | 300 | 0.230 (±0.000) | 653 | 0.234 (±0.000) | 680 | | `lang/statics_and_literals.hlsl` | 0.431 (±0.000) | 1184 | 0.216 (±0.001) | 725 | 0.442 (±0.001) | 1206 | 0.435 (±0.001) | 1210 | | `lang/templates.hlsl` | 0.586 (±0.012) | 1518 | 0.232 (±0.002) | 776 | 0.581 (±0.001) | 1526 | 0.585 (±0.001) | 1559 | "refused" marks a file the target does not support; it is left out rather than timed as a failure. ## Process time per entry point Median of 15 warm runs in milliseconds, with the fastest and slowest run in brackets, for each entry point. "vs DXC" is DXC's median divided by the OneShader median. | Entry point | DXC ms | NativeAOT ms | vs DXC | ReadyToRun ms | vs DXC | JIT ms | vs DXC | |---|--:|--:|---|--:|---|--:|---| | `2d/blit.hlsl:PSMain(ps)` | 6.40 (5.98–7.61) | 2.57 (2.33–3.88) | 2.49× faster | 42.2 (38.5–46.6) | 6.59× slower | 99.8 (98.2–106) | 15.6× slower | | `2d/blit.hlsl:VSMain(vs)` | 6.49 (6.02–7.83) | 2.56 (2.39–3.90) | 2.53× faster | 42.4 (40.0–47.2) | 6.52× slower | 101 (99.1–107) | 15.5× slower | | `2d/blur.hlsl:PSMain(ps)` | 7.32 (6.96–8.60) | 2.99 (2.52–5.00) | 2.45× faster | 49.0 (44.9–51.8) | 6.70× slower | 115 (113–123) | 15.8× slower | | `2d/blur.hlsl:VSMain(vs)` | 6.43 (6.13–7.75) | 2.72 (2.51–4.94) | 2.37× faster | 46.4 (41.0–48.7) | 7.22× slower | 108 (107–115) | 16.9× slower | | `2d/color_grade.hlsl:PSMain(ps)` | 7.43 (6.97–8.46) | 2.95 (2.54–5.47) | 2.52× faster | 48.8 (44.2–50.4) | 6.56× slower | 117 (113–127) | 15.7× slower | | `2d/color_grade.hlsl:VSMain(vs)` | 6.95 (6.23–9.36) | 2.71 (2.51–4.40) | 2.57× faster | 48.2 (42.6–50.5) | 6.94× slower | 112 (106–121) | 16.1× slower | | `2d/msdf_text.hlsl:PSMain(ps)` | 7.47 (7.04–8.57) | 3.12 (2.97–4.83) | 2.39× faster | 48.6 (43.9–51.3) | 6.51× slower | 114 (109–126) | 15.2× slower | | `2d/msdf_text.hlsl:VSMain(vs)` | 6.75 (6.22–8.14) | 2.78 (2.54–4.56) | 2.43× faster | 46.9 (42.5–49.3) | 6.94× slower | 109 (107–122) | 16.1× slower | | `2d/renderer/quad.hlsl:PSBlur(ps)` | 7.34 (6.85–8.52) | 3.27 (3.02–3.54) | 2.25× faster | 48.2 (44.6–52.9) | 6.56× slower | 116 (112–120) | 15.8× slower | | `2d/renderer/quad.hlsl:PSGrade(ps)` | 7.95 (7.55–8.95) | 3.42 (3.25–5.33) | 2.32× faster | 49.6 (43.2–52.8) | 6.23× slower | 117 (113–129) | 14.7× slower | | `2d/renderer/quad.hlsl:PSSprite(ps)` | 6.74 (6.30–8.99) | 3.04 (2.68–5.46) | 2.22× faster | 47.5 (44.3–50.8) | 7.05× slower | 111 (108–116) | 16.4× slower | | `2d/renderer/quad.hlsl:VSMain(vs)` | 7.76 (7.41–10.2) | 3.42 (3.03–5.55) | 2.27× faster | 49.0 (45.4–51.3) | 6.32× slower | 115 (114–121) | 14.9× slower | | `2d/renderer/tile.hlsl:PSFlat(ps)` | 8.15 (7.61–9.41) | 3.00 (2.71–3.76) | 2.72× faster | 48.5 (44.7–52.3) | 5.95× slower | 114 (110–118) | 13.9× slower | | `2d/renderer/tile.hlsl:PSIds(ps)` | 9.14 (8.47–10.4) | 3.25 (2.82–3.56) | 2.82× faster | 49.8 (45.2–52.0) | 5.45× slower | 118 (114–126) | 12.9× slower | | `2d/renderer/tile.hlsl:VSMain(vs)` | 9.53 (8.51–10.6) | 3.39 (3.12–4.66) | 2.81× faster | 48.5 (43.3–50.2) | 5.09× slower | 118 (115–126) | 12.3× slower | | `2d/renderer/ui.hlsl:PSMain(ps)` | 17.2 (15.7–18.7) | 6.45 (6.06–8.82) | 2.67× faster | 55.3 (50.2–57.1) | 3.21× slower | 125 (124–127) | 7.28× slower | | `2d/renderer/ui.hlsl:VSMain(vs)` | 9.30 (8.75–10.7) | 4.41 (3.89–6.69) | 2.11× faster | 50.2 (45.6–53.6) | 5.40× slower | 119 (116–121) | 12.8× slower | | `2d/sdf_text.hlsl:PSMain(ps)` | 7.49 (6.69–8.48) | 2.87 (2.67–4.68) | 2.61× faster | 48.6 (45.0–50.4) | 6.49× slower | 113 (110–116) | 15.1× slower | | `2d/sdf_text.hlsl:VSMain(vs)` | 6.80 (6.23–8.94) | 2.73 (2.52–3.49) | 2.49× faster | 46.5 (44.8–50.6) | 6.84× slower | 108 (106–114) | 15.9× slower | | `2d/sprite.hlsl:PSMain(ps)` | 6.36 (6.15–7.43) | 2.65 (2.41–4.38) | 2.40× faster | 47.3 (43.1–51.3) | 7.44× slower | 106 (105–117) | 16.7× slower | | `2d/sprite.hlsl:VSMain(vs)` | 6.48 (6.12–7.84) | 2.61 (2.31–5.55) | 2.48× faster | 45.7 (43.6–49.0) | 7.05× slower | 106 (105–116) | 16.4× slower | | `2d/tilemap.hlsl:PSMain(ps)` | 7.02 (6.74–9.96) | 2.87 (2.59–3.11) | 2.45× faster | 48.2 (45.5–51.3) | 6.86× slower | 112 (109–118) | 16.0× slower | | `2d/tilemap.hlsl:VSMain(vs)` | 7.62 (6.81–9.16) | 2.85 (2.71–4.82) | 2.67× faster | 47.1 (43.9–51.2) | 6.18× slower | 112 (109–119) | 14.6× slower | | `2d/ui_rect.hlsl:PSMain(ps)` | 9.95 (9.32–11.1) | 3.73 (3.36–5.78) | 2.67× faster | 50.9 (44.5–53.2) | 5.11× slower | 119 (114–124) | 12.0× slower | | `2d/ui_rect.hlsl:VSMain(vs)` | 7.03 (6.46–8.40) | 3.36 (3.02–3.65) | 2.09× faster | 49.0 (44.4–52.2) | 6.97× slower | 118 (111–127) | 16.8× slower | | `3d/bloom_downsample.hlsl:PSMain(ps)` | 11.0 (10.5–12.3) | 5.14 (4.43–6.03) | 2.14× faster | 52.3 (47.7–55.4) | 4.75× slower | 119 (113–131) | 10.8× slower | | `3d/bloom_downsample.hlsl:VSMain(vs)` | 7.17 (6.38–8.15) | 2.99 (2.62–4.94) | 2.40× faster | 47.8 (40.0–49.4) | 6.67× slower | 112 (106–119) | 15.6× slower | | `3d/bloom_upsample.hlsl:PSMain(ps)` | 7.70 (7.21–11.3) | 3.15 (2.83–5.36) | 2.44× faster | 46.4 (41.4–49.5) | 6.03× slower | 109 (107–116) | 14.2× slower | | `3d/bloom_upsample.hlsl:VSMain(vs)` | 6.51 (6.12–7.91) | 2.75 (2.54–4.39) | 2.37× faster | 45.2 (41.5–46.0) | 6.94× slower | 104 (103–114) | 16.0× slower | | `3d/clustered_lights.hlsl:PSMain(ps)` | 10.9 (9.93–13.1) | 4.64 (4.28–5.79) | 2.35× faster | 51.7 (47.9–53.8) | 4.73× slower | 122 (119–129) | 11.2× slower | | `3d/clustered_lights.hlsl:VSMain(vs)` | 7.39 (6.41–8.20) | 3.47 (3.13–3.85) | 2.13× faster | 48.7 (44.7–50.9) | 6.59× slower | 117 (112–128) | 15.8× slower | | `3d/ibl.hlsl:PSMain(ps)` | 8.31 (7.94–9.83) | 3.57 (3.27–5.80) | 2.33× faster | 49.8 (45.3–52.8) | 6.00× slower | 118 (116–129) | 14.2× slower | | `3d/ibl.hlsl:PSSky(ps)` | 8.44 (7.44–9.07) | 3.42 (3.13–6.33) | 2.46× faster | 48.8 (44.4–52.1) | 5.79× slower | 119 (115–123) | 14.1× slower | | `3d/ibl.hlsl:VSMain(vs)` | 7.77 (7.17–8.80) | 3.29 (2.92–3.75) | 2.36× faster | 48.6 (45.0–51.9) | 6.26× slower | 113 (111–119) | 14.6× slower | | `3d/ibl.hlsl:VSSky(vs)` | 7.61 (7.12–8.97) | 3.39 (3.03–6.16) | 2.24× faster | 48.4 (44.4–50.9) | 6.36× slower | 113 (111–117) | 14.9× slower | | `3d/pbr_forward.hlsl:PSMain(ps)` | 12.9 (12.2–14.1) | 5.47 (5.00–6.28) | 2.36× faster | 54.3 (51.7–56.7) | 4.22× slower | 127 (124–133) | 9.90× slower | | `3d/pbr_forward.hlsl:VSMain(vs)` | 8.32 (7.67–9.40) | 3.75 (3.54–6.05) | 2.22× faster | 49.9 (43.7–53.3) | 6.00× slower | 119 (116–130) | 14.3× slower | | `3d/shadow_cascades.hlsl:PSMain(ps)` | 12.5 (11.3–14.9) | 4.00 (3.83–4.81) | 3.14× faster | 51.4 (49.2–55.1) | 4.10× slower | 125 (122–131) | 9.98× slower | | `3d/shadow_cascades.hlsl:PSSpiral(ps)` | 12.9 (12.0–13.9) | 4.11 (3.61–4.79) | 3.13× faster | 51.6 (46.2–54.3) | 4.01× slower | 126 (123–132) | 9.76× slower | | `3d/shadow_cascades.hlsl:VSMain(vs)` | 8.19 (6.96–8.99) | 3.43 (3.08–5.52) | 2.39× faster | 47.7 (44.3–53.4) | 5.82× slower | 119 (115–125) | 14.5× slower | | `3d/shadow_depth.hlsl:VSMain(vs)` | 6.23 (5.63–8.00) | 2.61 (2.21–3.61) | 2.39× faster | 46.3 (43.1–48.4) | 7.44× slower | 107 (105–116) | 17.2× slower | | `3d/skinned.hlsl:PSMain(ps)` | 6.18 (5.85–8.11) | 2.82 (2.51–4.72) | 2.19× faster | 47.8 (43.0–50.8) | 7.74× slower | 109 (108–111) | 17.7× slower | | `3d/skinned.hlsl:VSMain(vs)` | 6.80 (6.59–9.80) | 3.03 (2.87–4.23) | 2.25× faster | 48.1 (45.7–49.2) | 7.07× slower | 110 (107–115) | 16.1× slower | | `3d/ssao.hlsl:PSMain(ps)` | 9.53 (9.11–10.8) | 3.64 (3.12–4.01) | 2.62× faster | 51.3 (48.9–53.9) | 5.38× slower | 121 (119–125) | 12.7× slower | | `3d/ssao.hlsl:VSMain(vs)` | 7.21 (6.78–8.53) | 3.04 (2.64–4.54) | 2.37× faster | 48.1 (44.3–50.7) | 6.67× slower | 115 (112–124) | 16.0× slower | | `3d/stage_io.hlsl:PSMain(ps)` | 7.93 (7.52–9.05) | 2.93 (2.68–3.59) | 2.70× faster | 48.4 (44.2–50.0) | 6.10× slower | 109 (108–133) | 13.8× slower | | `3d/stage_io.hlsl:VSMain(vs)` | 7.62 (7.19–8.79) | 3.03 (2.74–3.41) | 2.52× faster | 48.4 (44.2–51.9) | 6.36× slower | 108 (106–117) | 14.2× slower | | `3d/tonemap.hlsl:PSMain(ps)` | 8.24 (8.08–10.0) | 3.60 (3.09–4.41) | 2.29× faster | 50.0 (47.1–51.8) | 6.07× slower | 120 (119–129) | 14.6× slower | | `3d/tonemap.hlsl:VSMain(vs)` | 6.75 (6.34–7.49) | 2.93 (2.66–3.06) | 2.30× faster | 48.0 (44.7–50.2) | 7.11× slower | 114 (112–118) | 16.8× slower | | `compute/append_consume.hlsl:CSMain(cs)` | 6.60 (6.07–7.72) | 2.70 (2.44–2.96) | 2.44× faster | 45.9 (43.5–49.1) | 6.96× slower | 105 (103–115) | 15.9× slower | | `compute/cluster_build.hlsl:CSMain(cs)` | 7.68 (7.46–8.84) | 3.47 (3.11–4.12) | 2.21× faster | 49.6 (45.7–51.9) | 6.47× slower | 116 (114–125) | 15.1× slower | | `compute/cull_instances.hlsl:CSMain(cs)` | 8.74 (7.90–9.77) | 3.39 (3.05–5.06) | 2.58× faster | 49.7 (44.5–52.7) | 5.69× slower | 122 (116–125) | 13.9× slower | | `compute/depth_bounds.hlsl:CSMain(cs)` | 8.08 (7.16–8.64) | 3.03 (2.67–3.31) | 2.67× faster | 48.5 (46.4–52.2) | 6.00× slower | 118 (114–125) | 14.6× slower | | `compute/histogram.hlsl:CSMain(cs)` | 6.54 (6.01–7.13) | 2.80 (2.46–4.71) | 2.34× faster | 47.1 (44.8–48.6) | 7.20× slower | 107 (104–112) | 16.3× slower | | `compute/ibl_prefilter.hlsl:CSMain(cs)` | 11.1 (10.6–12.2) | 4.30 (4.03–5.57) | 2.59× faster | 51.9 (48.5–54.1) | 4.66× slower | 121 (119–129) | 10.9× slower | | `compute/image_filter.hlsl:CSMain(cs)` | 9.30 (8.48–10.2) | 3.31 (2.71–5.50) | 2.81× faster | 47.4 (44.9–49.2) | 5.10× slower | 113 (111–118) | 12.1× slower | | `compute/light_cull.hlsl:CSMain(cs)` | 7.66 (7.19–8.47) | 3.09 (2.76–5.33) | 2.48× faster | 48.2 (44.9–51.0) | 6.30× slower | 113 (111–119) | 14.8× slower | | `compute/particles.hlsl:CSMain(cs)` | 7.71 (7.37–9.05) | 3.12 (2.75–4.00) | 2.47× faster | 49.6 (45.1–53.1) | 6.43× slower | 114 (111–119) | 14.8× slower | | `compute/prefix_sum.hlsl:CSMain(cs)` | 6.58 (6.05–7.69) | 2.77 (2.57–3.11) | 2.38× faster | 46.4 (44.2–48.8) | 7.05× slower | 105 (102–111) | 15.9× slower | | `compute/stage_layouts.hlsl:CSMain(cs)` | 7.58 (7.25–9.14) | 3.03 (2.79–5.30) | 2.50× faster | 48.4 (44.0–50.8) | 6.39× slower | 109 (107–112) | 14.3× slower | | `lang/inheritance.hlsl:CSMain(cs)` | 7.42 (6.99–8.21) | 3.26 (3.03–4.24) | 2.28× faster | 48.6 (45.2–51.8) | 6.55× slower | 113 (109–118) | 15.3× slower | | `lang/members.hlsl:CSMain(cs)` | 7.25 (6.92–8.76) | 3.48 (2.85–3.94) | 2.08× faster | 48.0 (45.5–51.8) | 6.61× slower | 112 (111–121) | 15.5× slower | | `lang/operators.hlsl:CSMain(cs)` | 9.05 (8.54–10.3) | 4.32 (4.05–5.04) | 2.10× faster | 49.4 (46.6–53.2) | 5.46× slower | 112 (110–121) | 12.4× slower | | `lang/resource_arrays.hlsl:CSMain(cs)` | 8.24 (7.79–9.33) | 3.16 (2.82–3.98) | 2.61× faster | 46.8 (45.0–49.2) | 5.68× slower | 116 (112–127) | 14.0× slower | | `lang/resource_locals.hlsl:CSMain(cs)` | 8.36 (7.97–9.40) | 3.07 (2.92–5.03) | 2.72× faster | 47.6 (42.3–51.3) | 5.70× slower | 113 (110–116) | 13.5× slower | | `lang/sizeof_static.hlsl:CSMain(cs)` | 6.71 (6.38–8.42) | 3.05 (2.94–3.25) | 2.20× faster | 46.9 (43.0–47.9) | 7.00× slower | 110 (106–118) | 16.3× slower | | `lang/statics_and_literals.hlsl:CSMain(cs)` | 7.84 (7.58–9.43) | 3.55 (3.39–6.07) | 2.21× faster | 50.0 (46.6–53.5) | 6.37× slower | 119 (118–128) | 15.2× slower | | `lang/templates.hlsl:CSMain(cs)` | 8.95 (7.86–10.0) | 3.93 (3.77–5.75) | 2.28× faster | 49.6 (47.1–52.1) | 5.54× slower | 121 (119–133) | 13.5× slower | ### Cold runs The first run of each entry point after the corpus was copied to a fresh directory, one sample each (no spread to show). The page cache was not dropped (that needs privileges the benchmark does not use), so "cold" here means first run, not a cold disk: the JIT and ReadyToRun builds pay for loading and compiling code on every process start, and every process, warm or cold, is a new process. | Entry point | DXC ms | NativeAOT ms | vs DXC | ReadyToRun ms | vs DXC | JIT ms | vs DXC | |---|--:|--:|---|--:|---|--:|---| | `2d/blit.hlsl:PSMain(ps)` | 7.97 | 2.51 | 3.18× faster | 44.2 | 5.55× slower | 103 | 13.0× slower | | `2d/blit.hlsl:VSMain(vs)` | 7.78 | 2.72 | 2.86× faster | 42.9 | 5.52× slower | 99.5 | 12.8× slower | | `2d/blur.hlsl:PSMain(ps)` | 8.23 | 2.89 | 2.85× faster | 49.7 | 6.04× slower | 112 | 13.6× slower | | `2d/blur.hlsl:VSMain(vs)` | 6.68 | 2.72 | 2.46× faster | 45.8 | 6.85× slower | 109 | 16.3× slower | | `2d/color_grade.hlsl:PSMain(ps)` | 8.44 | 2.92 | 2.90× faster | 49.1 | 5.81× slower | 119 | 14.1× slower | | `2d/color_grade.hlsl:VSMain(vs)` | 7.29 | 2.77 | 2.63× faster | 45.1 | 6.19× slower | 114 | 15.7× slower | | `2d/msdf_text.hlsl:PSMain(ps)` | 9.07 | 3.06 | 2.97× faster | 48.0 | 5.29× slower | 121 | 13.3× slower | | `2d/msdf_text.hlsl:VSMain(vs)` | 8.54 | 2.67 | 3.19× faster | 46.5 | 5.45× slower | 110 | 12.9× slower | | `2d/renderer/quad.hlsl:PSBlur(ps)` | 8.30 | 2.93 | 2.83× faster | 50.4 | 6.07× slower | 112 | 13.5× slower | | `2d/renderer/quad.hlsl:PSGrade(ps)` | 8.27 | 3.99 | 2.07× faster | 46.0 | 5.57× slower | 124 | 15.0× slower | | `2d/renderer/quad.hlsl:PSSprite(ps)` | 7.97 | 2.88 | 2.76× faster | 45.5 | 5.71× slower | 113 | 14.2× slower | | `2d/renderer/quad.hlsl:VSMain(vs)` | 8.81 | 3.86 | 2.28× faster | 49.3 | 5.60× slower | 114 | 12.9× slower | | `2d/renderer/tile.hlsl:PSFlat(ps)` | 9.77 | 3.92 | 2.49× faster | 44.8 | 4.59× slower | 111 | 11.3× slower | | `2d/renderer/tile.hlsl:PSIds(ps)` | 10.6 | 4.14 | 2.56× faster | 48.1 | 4.53× slower | 115 | 10.9× slower | | `2d/renderer/tile.hlsl:VSMain(vs)` | 11.0 | 3.29 | 3.35× faster | 49.6 | 4.51× slower | 116 | 10.5× slower | | `2d/renderer/ui.hlsl:PSMain(ps)` | 18.8 | 6.56 | 2.87× faster | 52.4 | 2.78× slower | 127 | 6.73× slower | | `2d/renderer/ui.hlsl:VSMain(vs)` | 9.43 | 3.64 | 2.59× faster | 47.7 | 5.06× slower | 116 | 12.3× slower | | `2d/sdf_text.hlsl:PSMain(ps)` | 8.07 | 2.90 | 2.78× faster | 47.7 | 5.91× slower | 115 | 14.3× slower | | `2d/sdf_text.hlsl:VSMain(vs)` | 7.79 | 2.48 | 3.14× faster | 47.4 | 6.08× slower | 112 | 14.3× slower | | `2d/sprite.hlsl:PSMain(ps)` | 7.91 | 2.38 | 3.32× faster | 48.1 | 6.08× slower | 107 | 13.6× slower | | `2d/sprite.hlsl:VSMain(vs)` | 7.60 | 2.56 | 2.97× faster | 45.8 | 6.02× slower | 108 | 14.2× slower | | `2d/tilemap.hlsl:PSMain(ps)` | 8.41 | 2.58 | 3.26× faster | 48.6 | 5.79× slower | 110 | 13.1× slower | | `2d/tilemap.hlsl:VSMain(vs)` | 8.46 | 5.62 | 1.51× faster | 42.0 | 4.96× slower | 112 | 13.2× slower | | `2d/ui_rect.hlsl:PSMain(ps)` | 11.3 | 3.51 | 3.23× faster | 53.0 | 4.68× slower | 124 | 11.0× slower | | `2d/ui_rect.hlsl:VSMain(vs)` | 8.25 | 3.16 | 2.61× faster | 47.6 | 5.77× slower | 123 | 14.9× slower | | `3d/bloom_downsample.hlsl:PSMain(ps)` | 12.8 | 5.20 | 2.47× faster | 51.8 | 4.03× slower | 120 | 9.36× slower | | `3d/bloom_downsample.hlsl:VSMain(vs)` | 6.78 | 2.97 | 2.28× faster | 44.3 | 6.54× slower | 110 | 16.3× slower | | `3d/bloom_upsample.hlsl:PSMain(ps)` | 8.24 | 4.68 | 1.76× faster | 43.9 | 5.33× slower | 113 | 13.7× slower | | `3d/bloom_upsample.hlsl:VSMain(vs)` | 6.90 | 2.28 | 3.02× faster | 44.3 | 6.42× slower | 108 | 15.6× slower | | `3d/clustered_lights.hlsl:PSMain(ps)` | 12.8 | 4.34 | 2.95× faster | 50.1 | 3.91× slower | 122 | 9.54× slower | | `3d/clustered_lights.hlsl:VSMain(vs)` | 8.25 | 3.19 | 2.59× faster | 48.3 | 5.86× slower | 118 | 14.3× slower | | `3d/ibl.hlsl:PSMain(ps)` | 8.80 | 3.41 | 2.58× faster | 49.2 | 5.60× slower | 122 | 13.9× slower | | `3d/ibl.hlsl:PSSky(ps)` | 8.22 | 3.10 | 2.65× faster | 48.1 | 5.86× slower | 127 | 15.5× slower | | `3d/ibl.hlsl:VSMain(vs)` | 8.20 | 3.40 | 2.41× faster | 47.4 | 5.78× slower | 128 | 15.6× slower | | `3d/ibl.hlsl:VSSky(vs)` | 9.01 | 3.27 | 2.76× faster | 50.1 | 5.56× slower | 122 | 13.5× slower | | `3d/pbr_forward.hlsl:PSMain(ps)` | 16.1 | 7.26 | 2.22× faster | 50.0 | 3.10× slower | 131 | 8.14× slower | | `3d/pbr_forward.hlsl:VSMain(vs)` | 9.63 | 3.79 | 2.54× faster | 46.6 | 4.84× slower | 130 | 13.5× slower | | `3d/shadow_cascades.hlsl:PSMain(ps)` | 14.3 | 4.00 | 3.57× faster | 48.9 | 3.43× slower | 120 | 8.41× slower | | `3d/shadow_cascades.hlsl:PSSpiral(ps)` | 14.6 | 4.38 | 3.33× faster | 50.0 | 3.43× slower | 125 | 8.61× slower | | `3d/shadow_cascades.hlsl:VSMain(vs)` | 9.68 | 3.36 | 2.88× faster | 51.0 | 5.26× slower | 125 | 12.9× slower | | `3d/shadow_depth.hlsl:VSMain(vs)` | 7.89 | 2.51 | 3.14× faster | 43.5 | 5.50× slower | 107 | 13.6× slower | | `3d/skinned.hlsl:PSMain(ps)` | 6.73 | 2.62 | 2.57× faster | 47.1 | 7.00× slower | 107 | 15.9× slower | | `3d/skinned.hlsl:VSMain(vs)` | 7.47 | 2.93 | 2.54× faster | 46.0 | 6.16× slower | 109 | 14.6× slower | | `3d/ssao.hlsl:PSMain(ps)` | 10.2 | 4.36 | 2.34× faster | 46.9 | 4.59× slower | 126 | 12.3× slower | | `3d/ssao.hlsl:VSMain(vs)` | 9.45 | 2.88 | 3.28× faster | 47.6 | 5.04× slower | 114 | 12.1× slower | | `3d/stage_io.hlsl:PSMain(ps)` | 8.44 | 3.14 | 2.69× faster | 48.1 | 5.71× slower | 108 | 12.8× slower | | `3d/stage_io.hlsl:VSMain(vs)` | 7.77 | 2.79 | 2.78× faster | 47.3 | 6.09× slower | 113 | 14.6× slower | | `3d/tonemap.hlsl:PSMain(ps)` | 9.09 | 3.44 | 2.64× faster | 48.8 | 5.37× slower | 123 | 13.5× slower | | `3d/tonemap.hlsl:VSMain(vs)` | 8.03 | 2.90 | 2.77× faster | 45.6 | 5.67× slower | 113 | 14.1× slower | | `compute/append_consume.hlsl:CSMain(cs)` | 7.00 | 2.35 | 2.98× faster | 44.9 | 6.42× slower | 108 | 15.4× slower | | `compute/cluster_build.hlsl:CSMain(cs)` | 8.19 | 3.80 | 2.16× faster | 47.8 | 5.83× slower | 119 | 14.5× slower | | `compute/cull_instances.hlsl:CSMain(cs)` | 9.41 | 3.23 | 2.92× faster | 46.8 | 4.98× slower | 122 | 12.9× slower | | `compute/depth_bounds.hlsl:CSMain(cs)` | 7.61 | 2.93 | 2.60× faster | 47.6 | 6.25× slower | 115 | 15.1× slower | | `compute/histogram.hlsl:CSMain(cs)` | 6.47 | 3.04 | 2.13× faster | 44.9 | 6.95× slower | 104 | 16.1× slower | | `compute/ibl_prefilter.hlsl:CSMain(cs)` | 13.9 | 5.66 | 2.46× faster | 48.8 | 3.51× slower | 123 | 8.87× slower | | `compute/image_filter.hlsl:CSMain(cs)` | 10.6 | 3.14 | 3.36× faster | 47.4 | 4.49× slower | 114 | 10.8× slower | | `compute/light_cull.hlsl:CSMain(cs)` | 7.71 | 3.05 | 2.53× faster | 48.4 | 6.27× slower | 114 | 14.8× slower | | `compute/particles.hlsl:CSMain(cs)` | 9.64 | 3.26 | 2.96× faster | 45.6 | 4.73× slower | 126 | 13.0× slower | | `compute/prefix_sum.hlsl:CSMain(cs)` | 7.92 | 2.61 | 3.03× faster | 43.2 | 5.46× slower | 108 | 13.7× slower | | `compute/stage_layouts.hlsl:CSMain(cs)` | 7.79 | 2.76 | 2.82× faster | 47.7 | 6.12× slower | 110 | 14.1× slower | | `lang/inheritance.hlsl:CSMain(cs)` | 7.17 | 3.24 | 2.22× faster | 47.0 | 6.55× slower | 110 | 15.3× slower | | `lang/members.hlsl:CSMain(cs)` | 8.74 | 3.65 | 2.39× faster | 45.4 | 5.20× slower | 113 | 12.9× slower | | `lang/operators.hlsl:CSMain(cs)` | 10.1 | 4.34 | 2.32× faster | 48.2 | 4.79× slower | 116 | 11.5× slower | | `lang/resource_arrays.hlsl:CSMain(cs)` | 8.20 | 3.16 | 2.60× faster | 49.1 | 5.99× slower | 112 | 13.7× slower | | `lang/resource_locals.hlsl:CSMain(cs)` | 10.2 | 4.21 | 2.43× faster | 45.3 | 4.42× slower | 112 | 10.9× slower | | `lang/sizeof_static.hlsl:CSMain(cs)` | 7.56 | 2.99 | 2.53× faster | 47.9 | 6.34× slower | 108 | 14.3× slower | | `lang/statics_and_literals.hlsl:CSMain(cs)` | 8.90 | 3.36 | 2.65× faster | 48.7 | 5.47× slower | 132 | 14.8× slower | | `lang/templates.hlsl:CSMain(cs)` | 11.1 | 3.97 | 2.81× faster | 46.5 | 4.17× slower | 118 | 10.6× slower | ## Output size SPIR-V words (a word is 4 bytes) for the same entry points, from the NativeAOT build (the other two builds write identical bytes). "Without debug names" is the same compile with `KeepNames = false`. DXC is run with its defaults, which emit no names. | Entry point | DXC words | OneShader words | OneShader vs DXC | Without debug names | vs DXC | |---|--:|--:|--:|--:|--:| | `2d/blit.hlsl:PSMain(ps)` | 193 | 192 | -0.5% | 155 | -19.7% | | `2d/blit.hlsl:VSMain(vs)` | 187 | 202 | +8.0% | 178 | -4.8% | | `2d/blur.hlsl:PSMain(ps)` | 526 | 565 | +7.4% | 488 | -7.2% | | `2d/blur.hlsl:VSMain(vs)` | 187 | 202 | +8.0% | 178 | -4.8% | | `2d/color_grade.hlsl:PSMain(ps)` | 651 | 678 | +4.1% | 596 | -8.4% | | `2d/color_grade.hlsl:VSMain(vs)` | 187 | 202 | +8.0% | 178 | -4.8% | | `2d/msdf_text.hlsl:PSMain(ps)` | 650 | 676 | +4.0% | 594 | -8.6% | | `2d/msdf_text.hlsl:VSMain(vs)` | 237 | 246 | +3.8% | 194 | -18.1% | | `2d/renderer/quad.hlsl:PSBlur(ps)` | 466 | 496 | +6.4% | 442 | -5.2% | | `2d/renderer/quad.hlsl:PSGrade(ps)` | 680 | 673 | -1.0% | 618 | -9.1% | | `2d/renderer/quad.hlsl:PSSprite(ps)` | 221 | 214 | -3.2% | 174 | -21.3% | | `2d/renderer/quad.hlsl:VSMain(vs)` | 635 | 659 | +3.8% | 575 | -9.4% | | `2d/renderer/tile.hlsl:PSFlat(ps)` | 292 | 282 | -3.4% | 242 | -17.1% | | `2d/renderer/tile.hlsl:PSIds(ps)` | 622 | 599 | -3.7% | 529 | -15.0% | | `2d/renderer/tile.hlsl:VSMain(vs)` | 680 | 703 | +3.4% | 607 | -10.7% | | `2d/renderer/ui.hlsl:PSMain(ps)` | 2961 | 2770 | -6.5% | 2641 | -10.8% | | `2d/renderer/ui.hlsl:VSMain(vs)` | 846 | 838 | -0.9% | 690 | -18.4% | | `2d/sdf_text.hlsl:PSMain(ps)` | 555 | 581 | +4.7% | 498 | -10.3% | | `2d/sdf_text.hlsl:VSMain(vs)` | 237 | 246 | +3.8% | 194 | -18.1% | | `2d/sprite.hlsl:PSMain(ps)` | 265 | 267 | +0.8% | 223 | -15.8% | | `2d/sprite.hlsl:VSMain(vs)` | 278 | 287 | +3.2% | 223 | -19.8% | | `2d/tilemap.hlsl:PSMain(ps)` | 347 | 359 | +3.5% | 311 | -10.4% | | `2d/tilemap.hlsl:VSMain(vs)` | 517 | 549 | +6.2% | 454 | -12.2% | | `2d/ui_rect.hlsl:PSMain(ps)` | 1076 | 1118 | +3.9% | 947 | -12.0% | | `2d/ui_rect.hlsl:VSMain(vs)` | 612 | 747 | +22.1% | 554 | -9.5% | | `3d/bloom_downsample.hlsl:PSMain(ps)` | 1621 | 1746 | +7.7% | 1672 | +3.1% | | `3d/bloom_downsample.hlsl:VSMain(vs)` | 187 | 202 | +8.0% | 178 | -4.8% | | `3d/bloom_upsample.hlsl:PSMain(ps)` | 815 | 853 | +4.7% | 779 | -4.4% | | `3d/bloom_upsample.hlsl:VSMain(vs)` | 187 | 202 | +8.0% | 178 | -4.8% | | `3d/clustered_lights.hlsl:PSMain(ps)` | 1608 | 1720 | +7.0% | 1462 | -9.1% | | `3d/clustered_lights.hlsl:VSMain(vs)` | 456 | 470 | +3.1% | 333 | -27.0% | | `3d/ibl.hlsl:PSMain(ps)` | 827 | 853 | +3.1% | 710 | -14.1% | | `3d/ibl.hlsl:PSSky(ps)` | 543 | 553 | +1.8% | 432 | -20.4% | | `3d/ibl.hlsl:VSMain(vs)` | 382 | 396 | +3.7% | 270 | -29.3% | | `3d/ibl.hlsl:VSSky(vs)` | 187 | 202 | +8.0% | 178 | -4.8% | | `3d/pbr_forward.hlsl:PSMain(ps)` | 1953 | 2215 | +13.4% | 1881 | -3.7% | | `3d/pbr_forward.hlsl:VSMain(vs)` | 609 | 659 | +8.2% | 519 | -14.8% | | `3d/shadow_cascades.hlsl:PSMain(ps)` | 1428 | 1168 | -18.2% | 1009 | -29.3% | | `3d/shadow_cascades.hlsl:PSSpiral(ps)` | 1758 | 1385 | -21.2% | 1224 | -30.4% | | `3d/shadow_cascades.hlsl:VSMain(vs)` | 442 | 456 | +3.2% | 332 | -24.9% | | `3d/shadow_depth.hlsl:VSMain(vs)` | 348 | 380 | +9.2% | 301 | -13.5% | | `3d/skinned.hlsl:PSMain(ps)` | 396 | 425 | +7.3% | 359 | -9.3% | | `3d/skinned.hlsl:VSMain(vs)` | 728 | 1168 | +60.4% | 1055 | +44.9% | | `3d/ssao.hlsl:PSMain(ps)` | 1066 | 1159 | +8.7% | 1040 | -2.4% | | `3d/ssao.hlsl:VSMain(vs)` | 187 | 202 | +8.0% | 178 | -4.8% | | `3d/stage_io.hlsl:PSMain(ps)` | 587 | 620 | +5.6% | 539 | -8.2% | | `3d/stage_io.hlsl:VSMain(vs)` | 346 | 471 | +36.1% | 389 | +12.4% | | `3d/tonemap.hlsl:PSMain(ps)` | 926 | 1017 | +9.8% | 922 | -0.4% | | `3d/tonemap.hlsl:VSMain(vs)` | 187 | 202 | +8.0% | 178 | -4.8% | | `compute/append_consume.hlsl:CSMain(cs)` | 588 | 634 | +7.8% | 503 | -14.5% | | `compute/cluster_build.hlsl:CSMain(cs)` | 897 | 970 | +8.1% | 857 | -4.5% | | `compute/cull_instances.hlsl:CSMain(cs)` | 893 | 979 | +9.6% | 804 | -10.0% | | `compute/depth_bounds.hlsl:CSMain(cs)` | 631 | 627 | -0.6% | 547 | -13.3% | | `compute/histogram.hlsl:CSMain(cs)` | 478 | 514 | +7.5% | 425 | -11.1% | | `compute/ibl_prefilter.hlsl:CSMain(cs)` | 1249 | 1247 | -0.2% | 1160 | -7.1% | | `compute/image_filter.hlsl:CSMain(cs)` | 706 | 734 | +4.0% | 693 | -1.8% | | `compute/light_cull.hlsl:CSMain(cs)` | 808 | 923 | +14.2% | 697 | -13.7% | | `compute/particles.hlsl:CSMain(cs)` | 808 | 910 | +12.6% | 753 | -6.8% | | `compute/prefix_sum.hlsl:CSMain(cs)` | 451 | 475 | +5.3% | 398 | -11.8% | | `compute/stage_layouts.hlsl:CSMain(cs)` | 858 | 910 | +6.1% | 789 | -8.0% | | `lang/inheritance.hlsl:CSMain(cs)` | 406 | 486 | +19.7% | 379 | -6.7% | | `lang/members.hlsl:CSMain(cs)` | 487 | 609 | +25.1% | 525 | +7.8% | | `lang/operators.hlsl:CSMain(cs)` | 207 | 234 | +13.0% | 203 | -1.9% | | `lang/resource_arrays.hlsl:CSMain(cs)` | 806 | 855 | +6.1% | 775 | -3.8% | | `lang/resource_locals.hlsl:CSMain(cs)` | 551 | 609 | +10.5% | 540 | -2.0% | | `lang/sizeof_static.hlsl:CSMain(cs)` | 233 | 443 | +90.1% | 399 | +71.2% | | `lang/statics_and_literals.hlsl:CSMain(cs)` | 556 | 677 | +21.8% | 561 | +0.9% | | `lang/templates.hlsl:CSMain(cs)` | 631 | 709 | +12.4% | 649 | +2.9% | | **All 68 entry points** | **44126** | **46620** | **+5.7%** | **40258** | **-8.8%** | ## Method - **Entry points.** Every (file, entry point, stage) that the execution-parity cases in `corpus/cases/*.case` compile, without duplicates: 68 entry points in 41 files (a file's entry points are compiled together in the in-process benchmark and one per process in the process benchmark). Includes and the cases' preprocessor definitions are passed to both compilers. - **DXC.** The released Linux x64 binary v1.9.2609 (`libdxcompiler` 1.9.2609.5), pinned with its checksum in `tools/oracles.env`, run as `dxc -spirv -fspv-target-env=vulkan1.2 -HV 2021 -T _6_0 -E -I -Fo `; no optimization flags (so its default). It is used only as a black-box binary. - **OneShader.** `bench/OneShader.Bench.Cli` (`oshc`): `ShaderCompiler.CompileFile` with one entry point, SPIR-V, default options (optimizer on, debug names kept), writes the blob. Published three ways: `dotnet publish -r linux-x64 -p:PublishAot=true`; `--self-contained false -p:PublishReadyToRun=true`; and `--self-contained false` (JIT, with the runtime's default tiered compilation). The framework-dependent builds run through their `oshc` launcher on the machine's .NET runtime. - **Process timing.** The driver creates the process, waits for it to exit and takes the elapsed time on a monotonic clock; standard output and error are captured and discarded. Warm rounds run the tools interleaved (DXC, then each OneShader build) for each entry point, so a drift in machine speed affects all of them alike. Each round visits all 68 entry points; there are 15 rounds. - **In-process timing.** BenchmarkDotNet as described above. The in-process toolchain is used because the repository's build output layout does not suit BenchmarkDotNet's generated projects; the compiler runs in the benchmark's own process. - **Size.** Bytes written by each compiler on the cold pass, divided by 4. - **Reproduce.** `bench/run.sh` (optionally `--warm N`, `--skip-inproc`, `--skip-process`, `--out DIR`) builds, publishes the three builds, fetches the pinned DXC with `tools/fetch-oracles.sh`, runs both benchmarks, writes `bench/results/` and regenerates the tables on this page and the charts in `docs/img/`. It needs the .NET SDK from `global.json`, Python 3, and a C toolchain with zlib for NativeAOT. `bench/OneShader.Bench list` prints the 68 entry points. ### Machine and load - **Hardware and software.** AMD Ryzen 9 7950X3D (16 cores, 32 logical), 125 GiB RAM, Linux Mint 22.3, .NET SDK 10.0.402 with runtime 10.0.13, DXC v1.9.2609, BenchmarkDotNet 0.15.8. Default runtime settings (no AVX-512 switches, tiered compilation on, workstation non-concurrent GC in the benchmark process). - **A shared machine, run under a lock.** The run was the only benchmark on the machine: it held the machine's benchmark lock for the whole run and no other project's heavy job overlapped. Load averages (1 minute) were sampled every 5 seconds by the benchmark itself: at the start 1.24 (1-minute), 1.44 (5-minute), 1.49 (15-minute); during the process timing 1.03 to 1.24 over 41 samples; during the in-process benchmark 1.02 to 1.42 over 95 samples; at the end of the process timing 1.13, 1.26 and 1.41, and 1.09, 1.14 and 1.28 when the whole run finished. The remaining load of about 1 is background work on the machine that the benchmark did not control. - **One run.** The published numbers are from a single full run (2026-10-25). The spread between repeated warm runs is shown in the process table: a NativeAOT process lasts about 3 ms, so its slowest run is typically 1.5× its entry point's median (at most 2.1×), while DXC's slowest run is typically 1.15× its median (at most 1.5×). The ratios use medians. In-process results are kept as BenchmarkDotNet's `*-report-full.json` in `bench/results/inproc/`, with the file list per target (`params-*.tsv`) the report uses to restore the file names BenchmarkDotNet shortens. - **Not measured.** Memory use of either compiler (the NativeAOT binary is 2.9 MB; DXC is a 0.76 MB launcher plus a 31.8 MB compiler library); Windows or macOS; ARM64; multi-threaded or batch compiles (DXC and OneShader are each run one entry point at a time); compile time against other shader compilers. # Changelog All notable changes to the packages `RecursiveAI.OneShader`, `RecursiveAI.OneShader.Build` and `RecursiveAI.OneShader.Tool`. Versions follow [Semantic Versioning](https://semver.org/). ## [Unreleased] ## [1.2.0] - 16-bit and 64-bit types. With `CompileOptions.Enable16BitTypes` (DXC's `-enable-16bit-types`; MSBuild `OneShaderEnable16BitTypes`, item metadata `Enable16BitTypes`, CLI `-enable-16bit-types`) `half`, `min16float`, `min10float`, `min16int`, `min16uint` and `float16_t`/`int16_t`/`uint16_t` (vectors and matrices) are native 16-bit types, in arithmetic and in buffers with DXC's byte layout; without it the 16-bit names other than `half`/`min*` are error OSH0360, and the min-precision types are 32-bit values marked RelaxedPrecision on SPIR-V, as DXC does. `double`, `int64_t` and `uint64_t` work on SPIR-V in arithmetic, buffers, `asdouble`, `asuint(double, lo, hi)` and 64-bit `Interlocked*` on structured buffers and group-shared memory. A double argument to a float-only intrinsic (`floor`, `sqrt`, ...) is narrowed to float as DXC does, with the new warning OSH0361. Unsuffixed float literals keep double precision until they meet a concrete type. Constant folding is exact at each width. - MSL: native `half`, `short`, `ushort`, `long` and `ulong`; `double`, 64-bit integer vectors in buffers and 64-bit atomics other than unsigned `InterlockedMax`/`InterlockedMin` without a used result (the only form Metal 2.4 has) are errors at the first use (OSH5102) that say what to change. WGSL: `half` is `f16` with `enable f16;` (16-bit values in buffers are read and written as halves of 32-bit words; stores use atomic and/or on the word so neighbours written by other invocations are not disturbed); 16-bit integers, `double` and 64-bit integers are errors at the first use (OSH5202). A refusal is an error only when that target is selected. - Reflection: `EntryPointInfo.RequiredFeatures` (`DeviceFeatures`: `Float16`, `Int16`, `Float64`, `Int64`, `StorageBuffer16`, `UniformBuffer16`, `PushConstant16`, `InputOutput16`, `Int64Atomics`; the serialized `entry` record gets an optional `features` child). `EntryLayout.RequiredFeatures` and `PipelineLayout.RequiredFeatures`; `PipelineLayoutInfo.RequiredFeatures` (the Vulkan device features to enable); `WebGpuPipelineLayout.RequiredFeatures` lists `shader-f16`. Generated entry classes get a `RequiredFeatures` constant beside `SubgroupOperations`, and the generated `Layout` carries it. - Default vertex formats for 16-bit inputs: `Float16x2`/`Float16x4`, `Sint16x2`/`Sint16x4`, `Uint16x2`/`Uint16x4`. - Pipeline layout helpers (`OneShader.Layouts`), additions only. `PipelineLayout.TryCreateUnion` builds one layout for any number of entry points, several per stage; resources at one slot that agree merge with united stages, disagreements are a `LayoutError` naming both entry points and the slot, and push-constant blocks merge into one covering range when their overlapping members agree (`EntryLayout.PushConstantType`, `WithPushConstantType`). Entry points compiled from different files that disagree on a Metal index or WebGPU push slot are reported by `TryToMetal`/`TryToWebGpu`. `VulkanLayoutOptions` and `WebGpuLayoutOptions` gain stage widening (`WidenSets`/`WidenGroups`, `WidenBindings`, `WidenPushConstants`) and dynamic offsets by slot (`DynamicBindings`; WebGPU also `DynamicBuffers`); `ResourceLayout.DynamicOffset`, `DescriptorSetLayoutBindingInfo.IsDynamic` and `BindGroupLayoutEntryInfo.HasDynamicOffset` expose the flag. `EntryLayout.TryWithVertexOverrides` / `VertexLayout.TryWithOverrides` apply the host's stride, and per-location format and offset, checked against the shader's input type (`LayoutErrorKind.InvalidVertexInput`); Vulkan, Metal and WebGPU vertex descriptions follow. `TryToMetal` also reports two resources of one stage that take the same Metal index. - Specialization constants. `[[vk::constant_id(N)]] const Name = ;` becomes each target's pipeline-time constant, so one compiled shader serves many variants without recompiling. SPIR-V: `OpSpecConstant`/`OpSpecConstantTrue`/`False` with `SpecId N`, the ids DXC writes for the same source. MSL: `constant T Name [[function_constant(N)]]`, read through `is_function_constant_defined` with the default as the fallback. WGSL: `@id(N) override Name: T = default;`. Optimization never folds a specialization constant to its default; code derived from one (including `static const` initializers) is ordinary code that runs per pipeline. Mistakes are errors at the HLSL line: a duplicate id, a vector or struct type, `static`, a missing or non-constant default, a negative id (OSH0335); an array size, `[numthreads]` argument or `case` label that uses one says why it is refused (OSH0310), as DXC refuses the same uses. - Reflection: `ShaderReflection.SpecConstants` (`SpecConstantInfo`: name, id, scalar type, default bits and value, a `SpecConstantSlot` per target with the SPIR-V `SpecId`, the Metal function-constant index and variable, the WGSL `@id` and override name) and `EntryPointInfo.UsedSpecConstants`. The serialized form gets `spec` and `specuse` records, written only for shaders that have constants; the command-line tool's JSON gets `specConstants` and `usedSpecConstants` likewise. `InterfaceCompatibility` reports a removed constant or a changed id or type (`InterfaceDifferenceKind.SpecConstant`); new constants and changed defaults stay compatible. - `OneShader.Layouts`: `SpecializationValues` (from `EntryLayout.SpecConstants`; typed `Set`, plain data) converts to Vulkan `VkSpecializationInfo` contents (`Layouts.Vulkan.SpecializationInfo`: map entries and the data block), Metal function-constant values (`Layouts.Metal.FunctionConstantValue`: index, `MTLDataType` number, bits) and the WebGPU `constants` record keyed by decimal id. `EntryLayout` gets `SpecConstants`. - Generated code: `SpecializationConstants..Id`/`Default` constants, a blittable `Specialization` struct (one 4-byte field per constant, `Defaults`, `ToVulkan`/`ToMetal`/`ToWebGpu`) and `SpecConstantIds` on every entry-point class that reads constants; the entry's `Layout` carries them. In a variant file each variant class has its own `SpecializationConstants` and `Specialization` (a constant may exist in only some variants), and the aggregate entry class has `SpecConstantIds` for the variant it was looked up in. - `CompileOptions.SpecializationOverrides` fixes chosen constants at compile time: they become plain constants, get folded, and leave reflection. A value that does not fit the type is error OSH0335; an unknown name is warning OSH0335. - Metal note: a function that declares function constants cannot build a pipeline state as it is, even when every default is wanted; create it with `makeFunction(name:constantValues:)` (an empty `MTLFunctionConstantValues` selects the defaults). - Shader variants in the build package. A `OneShader` item takes `Axes="ALPHA_TEST;LIGHTING=NONE|LAMBERT|PBR"` (flags and choices; the cartesian product is compiled), or `OneShaderVariant` items name hand-picked define sets. One compile per variant per selected target, in parallel; only variants whose define set, shader, includes or options changed are recompiled. More than 64 variants is error OSH0905 (raise it with `MaxVariants` / `OneShaderMaxVariants`); malformed declarations are OSH0904 and OSH0906. An error that happens under some variants is reported once on the HLSL line with the define set (`(variant ALPHA_TEST=1, LIGHTING=PBR)`). - Generated code for a file with variants: `Key` (one property per axis, enums for choices), `Variants.Get(key)` returning a `Variant` with an `Entry` per entry point (`SpirV`, `Msl`, `Wgsl`, `Layout`, ...), `Bindings` and `Pipeline` per variant, and a static class per variant shaped like a single file's class. Constant-buffer, push-constant and vertex structs are generated once when all variants agree on the layout, per variant when they differ. Files without variants generate exactly the same C# as before. See `docs/generated-bindings.md`. - A file listed explicitly as a `OneShader` item (to give it metadata) is no longer added a second time by the default `**/*.hlsl` glob. - `ShaderCompileCache` (`RecursiveAI.OneShader`): thread-safe, bounded compile cache keyed by source, path, options and the content of the includes the compile read; concurrent calls for one key compile once; `Clear()` empties it. - The build task logs `OneShader: N compiled, M up to date.` at low importance (`-v:d`). ## [1.1.0] - New package `RecursiveAI.OneShader.Tool`: the `oneshader` command-line compiler (a .NET tool, and a NativeAOT single-file binary with `dotnet publish -r `). It takes DXC's core flags (`-T`, `-E`, `-D`, `-I`, `-Fo`, `-spirv`, `-O0`..`-O3`, `-Zpr`/`-Zpc`, `-HV 2021`, `-WX`) plus `--target spirv|msl|wgsl` (repeatable), `--all-entries`, `--reflect ` and `--depfile `. Diagnostics print as `file(line,col): error OSHnnnn: message`; exit code 0, 1 for compile errors, 2 for command-line errors. - Pipeline layout helpers (`OneShader.Layouts`): plain data, no graphics API behind it, that turns a shader into what a renderer creates. `EntryLayout` (from reflection with `EntryLayout.TryCreate`, or the generated `Sprite.VSMain.Layout`) and `PipelineLayout.TryCreate` merge entry points into bind groups with combined stage visibility and push-constant blocks, returning conflicts as a `LayoutError` value. `ToVulkan()` gives descriptor set layouts, push-constant ranges and vertex input with numeric `VkDescriptorType`/`VkShaderStageFlagBits`/`VkFormat` values; `TryToMetal()` gives per-stage buffer, texture and sampler indices, the `setBytes` slot and `MTLVertexDescriptor` data; `TryToWebGpu()` gives bind group layout entries with the WebGPU enum strings, push constants as a uniform buffer at the reserved group and binding, and vertex buffer layouts. Generated classes gain `Layout` on every entry point and `Pipeline` (vertex + pixel, or one compute entry point) on the file. - Wave (subgroup) and quad intrinsics on all three targets: `WaveGetLaneCount`/`Index`, `WaveIsFirstLane`, `WaveActiveAnyTrue`/`AllTrue`/`AllEqual`/`Ballot`/`Sum`/`Product`/`Min`/`Max`/`BitAnd`/`BitOr`/`BitXor`/`CountBits`, `WavePrefixSum`/`Product`/`CountBits`, `WaveReadLaneAt`/`First` and `QuadReadAcrossX`/`Y`/`Diagonal`/`QuadReadLaneAt` on 32-bit int, uint and float scalars and vectors. SPIR-V uses `GroupNonUniform*` as DXC does, MSL `simd_*`/`quad_*`, WGSL the `subgroups` feature (`enable subgroups;`; not in vertex shaders). `WaveMatch` and `WaveMultiPrefix*` are an error (OSH0350). Reflection gains `EntryPointInfo.SubgroupOperations` (Vulkan's `VkSubgroupFeatureFlagBits` values), `EntryPointInfo.UsesSubgroups` and `ShaderReflection.UsesSubgroups`; `EntryLayout`/`PipelineLayout`/`PipelineLayoutInfo`/`MetalPipelineLayout` carry `SubgroupOperations`, `WebGpuPipelineLayout.RequiredFeatures` lists `subgroups`, and generated entry classes get a `SubgroupOperations` constant. `BuiltIn` gains `SubgroupLocalInvocationId` and `SubgroupSize`. - Reflection: `ResourceInfo.ImageFormat` records the `[[vk::image_format]]` name of a storage texture (the serialized `resource` record gets an optional `format` attribute, written only when set). ## [1.0.0] First stable release. The public API (`ShaderCompiler`, `CompileOptions`, `CompileResult`, reflection, `InterfaceCompatibility`) and the generated C# shape are now covered by Semantic Versioning. - Vertex, pixel and compute shaders in modern HLSL (HLSL 2021 templates included) compile to optimized SPIR-V for Vulkan 1.2, Metal Shading Language and WGSL, each with reflection; the build package compiles `*.hlsl` during `dotnet build`, reports errors on the HLSL line, embeds the selected targets (`OneShaderTargets`) and generates C# structs and binding code. - Since 0.4: the packages are built with deterministic paths, embedded untracked sources and a symbol package (`.snupkg`); package metadata links to the public repository; the four-platform CI (Windows x64, Linux x64 and arm64, macOS arm64) runs the compiler tests, the build-package tests and the NativeAOT smoke app, and runs the reference-tool tests (DXC, `spirv-val`, naga) wherever those tools are released. - Tested: the five fuzz targets (preprocessor, front end, IR and optimizer, every backend) ran clean for 600 s each, and the corpus is checked with `spirv-val`, by running it on GPUs against DXC's output, with naga, with Apple's `metal` compiler and on WebGPU. ## [0.4.0-preview.1] - Hardening from fuzzing every compile path: crashes, hangs and invalid SPIR-V or WGSL on unusual input are now diagnostics. Errors are capped at 100 per compile (OSH0004) and identical consecutive errors are dropped; array dimensions are limited to 16, local array initializers to 4096 elements, and binding, set and register numbers to 65535; the preprocessor's default output and expansion limits are 2M tokens and 20M steps. Component-wise intrinsics on matrices (`saturate`, `lerp`, `pow`, `sin` and others) work on every target. - Very deep call chains no longer overflow the stack on threads with a small stack (seen on macOS). - WGSL: `WgslBackend.EntryFunctionName` gives the name an entry point has in the WGSL text (reserved words such as `loop` or `fragment` get a trailing underscore). The build task uses it for `WgslEntryPoint` instead of reading the generated text, and a test checks entry points named after reserved words in MSL and WGSL against naga and `metal`. - MSL: `switch` cases that fall through and a `default` in the middle now run correctly. Before, a case entered by falling through re-ran the switch edge's phi copies and overwrote the value it fell in with, and a `break` or `continue` inside an `if` inside a case was dropped. WGSL, which has no fall-through, repeats the following case's statements; both are covered by the new `lang_switch_flow` parity case. - MSL honours `precise` under Metal's default fast math: functions with precise arithmetic turn off `reassociate` and `contract`, and precise divisions use `precise::divide`. - `precise` on local variables, assignments and function return values now marks the arithmetic that feeds them in every target (the SPIR-V output gets `NoContraction` on the same operations as DXC; before, the modifier was parsed and ignored). - WGSL: arithmetic marked `precise` compiles with warning OSH5205, because WGSL cannot forbid a fused multiply-add. ## [0.3.0-preview.1] - Push constants: `[[vk::push_constant]]` on a struct-typed global or a `ConstantBuffer` compiles on every target. SPIR-V gets a `PushConstant` block whose member offsets equal DXC's (vector-relaxed std430, not the cbuffer rules); MSL takes it as a `constant T&` at a Metal buffer index (set it with `setVertexBytes`, `setFragmentBytes` or `setBytes`); WGSL binds it as a uniform buffer at the next free binding of the highest bind group (WebGPU has no push constants). Reflection gains `ShaderReflection.PushConstants` (layout, stages, per-target slots) and `EntryPointInfo.PushConstant`; the generated class gains a `PushConstants` struct (offsets and `Size` equal to SPIR-V) and `Bindings.PushConstants` (`Size`, `Stages` for the `VkPushConstantRange`, `MslIndex`, `WgslGroup`/`WgslBinding`). An entry point that reads two blocks is OSH0333 at the second block; a block over 128 bytes (Vulkan's guaranteed `maxPushConstantsSize`) is warning OSH0334; `[[vk::push_constant]]` on a `cbuffer` is OSH0333 (DXC rejects it too). - Build package: `OneShaderTargets` (`SpirV;Msl;Wgsl`, default `SpirV`; per-item `Targets` metadata) embeds Metal and WGSL sources next to the SPIR-V. Entry-point classes gain `Msl`, `MslFunctionName`, `MslVertexBufferIndex`, `Wgsl` and `WgslEntryPoint`; resource bindings gain `MslIndex` / `MslCounterIndex` and `ResourceBinding.MetalIndex`. A target's refusal (for example OSH5202 on WGSL) is a build error on the HLSL line only when that target is selected; unknown target names are OSH0903. - New WGSL target (`CompileTarget.Wgsl`): vertex, fragment and compute entry points as WGSL text for WebGPU, with the same buffer byte layouts as the SPIR-V target. See docs/wgsl-backend.md for what is not available. - New: Metal Shading Language output (`CompileTarget.Msl`, `MslBackend`): one source file per entry point with the same buffer byte layout as the SPIR-V output, Metal argument indices in reflection, and `tools/validate-msl.sh` to check the files with Apple's compiler. See `docs/msl-backend.md`. ## [0.2.0-preview.1] - Optimizer, on by default (`CompileOptions.Optimize`): inlining, SSA promotion, constant folding, common-subexpression and dead-code elimination, control-flow simplification. SPIR-V is about 38% smaller than before and close to DXC's size. - HLSL 2021 language features: function and struct templates, member functions and `this`, single inheritance, operator overloads, `sizeof`, static locals, resource-typed locals and parameters. - Bounded resource arrays end to end (`Texture2D T[4]`, `SamplerState S[2]`, arrays of structured buffers and constant buffers): constant and run-time indexing, `NonUniformResourceIndex`, elements passed to functions, reflection and generated `Count`. Unbounded arrays are a clear error. - `AppendStructuredBuffer` and `ConsumeStructuredBuffer` (`Append`, `Consume`, `GetDimensions`) with the hidden counter exposed in reflection (`counter` records) and in the generated bindings (`CounterBinding`, a `.counter` entry in `All`). - A constant buffer may share its name with one of its members, and a whole `ConstantBuffer` can be copied into a struct local (`S s = cb;`). - Static struct members, member functions defined outside their struct, template argument deduction through struct templates (`Box`), and integer literal arithmetic in 64 bits (`(1 << 31) >> 3` is `0x10000000`, as in DXC). - A resource array occupies one binding number (it used to take one per element when numbering unannotated resources). - Stage inputs and outputs may be arrays and matrices (one location per element or row, like DXC), and `SV_ClipDistance` / `SV_CullDistance` work in vertex outputs and pixel inputs, as scalars, vectors or arrays. - `StructuredBuffer` and other elements whose size is not a multiple of their alignment use the aligned stride (16 bytes for `float3`), which Vulkan validation requires; `bool` members of constant and structured buffers are stored as 32-bit integers. - New: `Texture2D.CalculateLevelOfDetail` that validates, `GatherCmp`/`GatherCmpRed` on 2D, array and cube textures, `f16tof32`, `f32tof16`, `faceforward`, `log10`. - Overlapping `packoffset` members are an error (OSH0332) instead of invalid SPIR-V. - `ShaderCompiler.Compile` turns an unexpected internal exception into error OSH7000 on line 1 of the file instead of throwing; the build task reports internal errors at line 1 as well. - Hot reload support: `CompileResult.Files` lists the files a compile read (for file watchers), `ShaderCompiler.CompileFile` compiles from disk and reports an unreadable file as diagnostic OSH0003, and `InterfaceCompatibility.Compare` tells whether a recompiled shader still matches the generated C# (constant-buffer offsets and types, vertex inputs, bindings, entry points, thread-group sizes), listing each difference. Generated shader classes gain a `Reflection` member with the reflection they were generated from. Concurrent `Compile` calls are tested to give the same bytes as serial ones. `samples/HotReload` shows the loop. ## [0.1.0-preview.1] First preview. - HLSL compiler for vertex, pixel and compute shaders (modern HLSL, HLSL 2021): preprocessor, constant buffers, textures and samplers, structured and byte-address buffers, storage images, groupshared memory, barriers and atomics. - SPIR-V output for Vulkan 1.2, with reflection; buffer layouts and descriptor bindings follow DXC's Vulkan conventions. - MSBuild integration: `*.hlsl` compiled during `dotnet build`, errors on the HLSL line, embedded SPIR-V and generated C# structs, vertex inputs and binding constants. - Known limits: SPIR-V only; single resources only (no resource arrays); no templates, member functions, operator overloads or wave intrinsics yet; `dotnet build` only. [1.2.0]: https://github.com/recursive-ai-oss/oneshader-csharp/releases/tag/v1.2.0 [1.1.0]: https://github.com/recursive-ai-oss/oneshader-csharp/releases/tag/v1.1.0 [1.0.0]: https://github.com/recursive-ai-oss/oneshader-csharp/releases/tag/v1.0.0 [0.4.0-preview.1]: https://github.com/recursive-ai-oss/oneshader-csharp/releases/tag/v0.4.0-preview.1 [0.3.0-preview.1]: https://github.com/recursive-ai-oss/oneshader-csharp/releases/tag/v0.3.0-preview.1 [0.2.0-preview.1]: https://github.com/recursive-ai-oss/oneshader-csharp/releases/tag/v0.2.0-preview.1 [0.1.0-preview.1]: https://github.com/recursive-ai-oss/oneshader-csharp/releases/tag/v0.1.0-preview.1 # Compile HLSL to SPIR-V for Vulkan Compile an HLSL vertex or pixel shader to SPIR-V for Vulkan with the `oneshader` command and its DXC-style flags, or from C# with `ShaderCompiler`, and check the result with `spirv-val`. OneShader writes SPIR-V 1.5 for the Vulkan 1.2 environment, with buffer layouts and descriptor bindings that follow DXC’s Vulkan rules. The sample shader is `Sprite.hlsl` from [Getting started](https://recursiveai.org/oneshader/docs/getting-started/): a constant buffer at `b0`, a texture at `[[vk::binding(1)]]` and a sampler at `[[vk::binding(2)]]`. ## From the command line ```sh dotnet tool install -g RecursiveAI.OneShader.Tool --version 1.2.0 oneshader -T vs_6_0 -E VSMain -spirv Sprite.hlsl -Fo sprite.vs.spv oneshader -T ps_6_0 -E PSMain -spirv Sprite.hlsl -Fo sprite.ps.spv spirv-val --target-env vulkan1.2 sprite.vs.spv spirv-val --target-env vulkan1.2 sprite.ps.spv spirv-dis sprite.ps.spv | grep -E "Binding|DescriptorSet" ``` Both files pass `spirv-val` (from the Vulkan SDK) silently, and the bindings are where the shader put them: ```text OpDecorate %Atlas DescriptorSet 0 OpDecorate %Atlas Binding 1 OpDecorate %AtlasSampler DescriptorSet 0 OpDecorate %AtlasSampler Binding 2 ``` The same command line with `dxc` instead of `oneshader` gives the same set and binding numbers. ## From C# With the `RecursiveAI.OneShader` package (`dotnet add package RecursiveAI.OneShader --version 1.2.0`): ```csharp using OneShader; string source = File.ReadAllText("Sprite.hlsl"); CompileResult result = ShaderCompiler.Compile(source, "Sprite.hlsl", new CompileOptions { // The C# form of -T ps_6_0 -E PSMain; leave EntryPoints empty to compile every [shader("...")] function. EntryPoints = [new EntryPointRequest("PSMain", ShaderStage.Pixel)], }); foreach (var d in result.Diagnostics) Console.WriteLine(d); if (!result.Success) return 1; ShaderBlob ps = result.Find("PSMain", CompileTarget.SpirV)!; File.WriteAllBytes("sprite.ps.spv", ps.ToArray()); uint magic = BitConverter.ToUInt32(ps.Data[..4]); Console.WriteLine($"{ps.EntryPoint}: {ps.Data.Length} bytes, magic 0x{magic:x8}"); return 0; ``` ```text PSMain: 856 bytes, magic 0x07230203 ``` `0x07230203` is the SPIR-V magic number. The file written here is byte for byte what the command line wrote, and it passes `spirv-val` too. ## Good to know - **Bindings.** Use `[[vk::binding(n)]]` or `[[vk::binding(n, set)]]`, or HLSL registers, which map to bindings as DXC maps them by default. Unannotated resources are numbered in declaration order. `-fvk-*` remapping options aren’t supported: put the numbers in the shader. - **Debug names.** By default the SPIR-V keeps source names (`%Atlas` above). `--no-names` (or `KeepNames = false`) leaves them out: here the pixel shader goes from 856 to 692 bytes. - **Optimization** is on by default (`-O3` and `-O0` work as in DXC; `Optimize = false` in C#). - **Validation.** The repository validates every corpus shader with `spirv-val` and runs it on GPUs next to the same shader compiled by DXC, with the Vulkan validation layers on. - **Not in 1.2.0:** DXIL (Direct3D 12) output, and geometry, tessellation, mesh and ray-tracing stages. `-T gs_6_0` is a command-line error naming the supported profiles. The full SPIR-V mapping is in [SPIR-V (Vulkan)](https://recursiveai.org/oneshader/docs/spirv-backend/). To compile the same shader during `dotnet build`, see [Compile HLSL in a C# project](https://recursiveai.org/oneshader/cookbook/compile-hlsl-csharp/). Choosing a compiler: [HLSL to SPIR-V compiler: DXC, glslang and OneShader compared](https://recursiveai.org/oneshader/articles/hlsl-to-spirv-compiler/). # Compile HLSL in a C# project with dotnet build Compile HLSL shaders as part of `dotnet build` in a C# project, with shader errors reported as build errors on the HLSL line, shared include files, defines, a custom namespace and per-file targets. Reference `RecursiveAI.OneShader.Build` and every `*.hlsl` file in the project is compiled on build into SPIR-V embedded in the assembly, next to generated C# for its constant buffers, bindings and entry points. This recipe shows the project-level settings. [Getting started](https://recursiveai.org/oneshader/docs/getting-started/) shows the generated code itself. ## The project ```xml Exe net10.0 enable enable Game.Rendering.Shaders Shaders/Common RADIUS=4 ``` `Shaders/Common/Blur.hlsli`, shared by two variants: ```hlsl Texture2D Source : register(t0); SamplerState Linear : register(s1); cbuffer Blur : register(b2) { float2 TexelSize; }; [shader("pixel")] float4 PSMain(float2 uv : TEXCOORD0) : SV_Target { float2 step = HORIZONTAL ? float2(TexelSize.x, 0) : float2(0, TexelSize.y); float4 sum = 0; [unroll] for (int i = -RADIUS; i <= RADIUS; i++) sum += Source.Sample(Linear, uv + step * i); return sum / (2 * RADIUS + 1); } ``` `Shaders/BlurH.hlsl`: ```hlsl #define HORIZONTAL 1 #include "Blur.hlsli" ``` `Shaders/BlurV.hlsl`: ```hlsl #define HORIZONTAL 0 #include "Blur.hlsli" ``` `Shaders/Present.hlsl`, a full-screen triangle: ```hlsl Texture2D Frame : register(t0); SamplerState Point : register(s1); struct VSOut { float4 Position : SV_Position; float2 Uv : TEXCOORD0; }; [shader("vertex")] VSOut VSMain(uint id : SV_VertexID) { VSOut o; o.Uv = float2((id << 1) & 2, id & 2); o.Position = float4(o.Uv * float2(2, -2) + float2(-1, 1), 0, 1); return o; } [shader("pixel")] float4 PSMain(VSOut i) : SV_Target { return Frame.Sample(Point, i.Uv); } ``` `Program.cs`: ```csharp using Game.Rendering.Shaders; Console.WriteLine($"BlurH.PSMain: {BlurH.PSMain.SpirV.Length} bytes of SPIR-V"); Console.WriteLine($"BlurV.PSMain: {BlurV.PSMain.SpirV.Length} bytes of SPIR-V"); Console.WriteLine($"Blur constant buffer: {BlurH.Blur.Size} bytes at binding {BlurH.Bindings.Blur.Binding}"); Console.WriteLine($"Present.VSMain: {Present.VSMain.SpirV.Length} bytes of SPIR-V, {Present.VSMain.Wgsl.Length} bytes of WGSL"); ``` ```text BlurH.PSMain: 1580 bytes of SPIR-V BlurV.PSMain: 1580 bytes of SPIR-V Blur constant buffer: 16 bytes at binding 2 Present.VSMain: 836 bytes of SPIR-V, 740 bytes of WGSL ``` ## An error in an included file Misspell `step` in `Blur.hlsli` and `dotnet build` stops on the include’s own line: ```text /tmp/oneshader-docs/cookbook/Build/Shaders/Common/Blur.hlsli(11,43): error OSH0300: 'stepp' is not declared. Did you mean 'step'? [/tmp/oneshader-docs/cookbook/Build/Build.csproj] ``` ## Good to know - **Default items.** Every `*.hlsl` in the project is a `OneShader` item. `.hlsli` files aren’t, so they’re only compiled where they’re included. To give one file metadata, list it with ``, as above: since 1.2.0 the default glob skips files that are already listed. `` doesn’t work, because the package adds the default items after your project file. Setting `EnableDefaultOneShaderItems` to `false` turns the glob off. - **One class per file.** The class is named after the file in PascalCase, in `$(OneShaderNamespace)` (default `.Shaders`). Two `OneShader` items for the same file with different `ClassName` and `Defines` give only one class, without a warning. For variants of one source, use `Axes` or `OneShaderVariant` items ([Shader variants with Axes](https://recursiveai.org/oneshader/cookbook/shader-variants-axes/)), or give each variant its own small `.hlsl` file, as above. - **Item metadata**, all optional: `EntryPoints` (`VSMain:Vertex;PSMain:Pixel`) (instead of the `[shader]` attributes), `Defines="A;B=2"`, `IncludeDirectories`, `Namespace`, `ClassName`, `Targets`, `Enable16BitTypes`, `Axes`, `MaxVariants`. **Project properties:** `OneShaderNamespace`, `OneShaderTargets` (default `SpirV`), `OneShaderDefines`, `OneShaderIncludeDirectories`, `OneShaderEnable16BitTypes`, `OneShaderMaxVariants`. - **Includes** are searched in the including file’s folder first, then in the include directories. - **Name clashes are errors.** A `cbuffer` named like its file (`cbuffer Lighting` in `Lighting.hlsl`) is error OSH0951, because the struct and the class would share a name. - **Incremental.** A shader is compiled again only when it, a file it includes, the options or the compiler change. The output is deterministic. `dotnet build -v:d` logs the count; after the typo above was fixed it was `OneShader: 2 compiled, 1 up to date.` (the two blur files share the include). - **The build package doesn’t flow** to projects that reference yours: only projects with shaders reference it. The support types live in `RecursiveAI.OneShader`, so a rendering layer in another assembly can take them. - **`dotnet build` only.** Visual Studio’s .NET Framework MSBuild can’t load the task. Every generated member is listed in [Generated bindings](https://recursiveai.org/oneshader/docs/generated-bindings/). For variants compiled at run time instead, see [Shader variants with defines and includes](https://recursiveai.org/oneshader/cookbook/shader-variants-defines-includes/). The options compared, dxc, Vortice.Dxc and OneShader: [Compile HLSL shaders in C#](https://recursiveai.org/oneshader/articles/csharp-shader-compiler/). # Match HLSL constant buffers with C# structs Get a C# struct whose field offsets equal an HLSL constant buffer’s packing, 16-byte rows, arrays and matrices included, so you can upload its bytes without hand-written padding. HLSL packs constant buffers in 16-byte rows. A vector can’t cross a row boundary, and each array element starts a new row. A hand-written C# struct has to repeat those rules exactly, and a mistake shows up as wrong values on screen. With `RecursiveAI.OneShader.Build`, the struct is generated from the compiled shader on every build, so it can’t drift from it. `Shaders/Lighting.hlsl`: ```hlsl cbuffer Light : register(b0) { float3 LightDirection; // 12 bytes, and Intensity fits in the same 16-byte row float Intensity; float2 ShadowBias; float3 Ambient; // would cross a 16-byte row after ShadowBias, so it starts the next one float Weights[4]; // each element of an array starts a new row float4x4 LightViewProjection; }; [shader("pixel")] float4 PSMain(float3 normal : NORMAL) : SV_Target { float l = saturate(dot(normal, -LightDirection)) * Intensity; float w = Weights[0] + Weights[1] + Weights[2] + Weights[3]; return float4(Ambient + l * w + ShadowBias.x, LightViewProjection[0][0]); } ``` `Program.cs`, in a project that references `RecursiveAI.OneShader.Build` 1.2.0: ```csharp using System.Numerics; using System.Reflection; using System.Runtime.InteropServices; using Packing.Shaders; // The generated struct has the shader's offsets, so its bytes can be uploaded as they are. Console.WriteLine($"Light.Size = {Lighting.Light.Size}, Marshal.SizeOf = {Marshal.SizeOf()}"); foreach (FieldInfo field in typeof(Lighting.Light).GetFields(BindingFlags.Instance | BindingFlags.Public | BindingFlags.NonPublic)) Console.WriteLine($"{field.Name,-24} {field.FieldType.Name,-8} offset {Marshal.OffsetOf(field.Name)}"); var light = new Lighting.Light { LightDirection = Vector3.Normalize(new Vector3(0, -1, -1)), Intensity = 2.5f, Weights_0 = 0.25f, Weights_1 = 0.25f, Weights_2 = 0.25f, Weights_3 = 0.25f, LightViewProjection = Matrix4x4.Identity, // element (i, j) is HLSL's LightViewProjection[i][j] }; ReadOnlySpan bytes = MemoryMarshal.AsBytes(new ReadOnlySpan(in light)); Console.WriteLine($"{bytes.Length} bytes to upload"); ``` ```text Light.Size = 176, Marshal.SizeOf = 176 LightDirection Vector3 offset 0 Intensity Single offset 12 ShadowBias Vector2 offset 16 Ambient Vector3 offset 32 Weights_0 Single offset 48 Weights_1 Single offset 64 Weights_2 Single offset 80 Weights_3 Single offset 96 _LightViewProjection_c0 Vector4 offset 112 _LightViewProjection_c1 Vector4 offset 128 _LightViewProjection_c2 Vector4 offset 144 _LightViewProjection_c3 Vector4 offset 160 176 bytes to upload ``` DXC agrees. Its SPIR-V for the same file (`dxc -T ps_6_0 -E PSMain -spirv`) declares: ```text OpDecorate %_arr_float_uint_4 ArrayStride 16 OpMemberDecorate %type_Light 0 Offset 0 OpMemberDecorate %type_Light 1 Offset 12 OpMemberDecorate %type_Light 2 Offset 16 OpMemberDecorate %type_Light 3 Offset 32 OpMemberDecorate %type_Light 4 Offset 48 OpMemberDecorate %type_Light 5 Offset 112 ``` ## What the generator does - **Explicit layout.** The struct is `[StructLayout(LayoutKind.Explicit)]` with a `FieldOffset` per member and a `uint Size` constant. It’s blittable, so `MemoryMarshal.AsBytes` gives the bytes to copy into the buffer. - **Arrays of scalars** become one field per element (`Weights_0` to `Weights_3`), each at the start of its own row, the 16-byte stride HLSL uses. - **Matrices** are a `System.Numerics.Matrix4x4` property over four private column vectors, because HLSL’s default packing is column-major. Set it as a normal `Matrix4x4`: element (i, j) is HLSL’s `M[i][j]`. A `row_major` matrix is stored the way `Matrix4x4` is, so it becomes a plain `Matrix4x4` field. Either way the C# code is the same. - **The same bytes for every target.** SPIR-V, Metal and WGSL use one byte layout, so one struct serves all three. The MSL output carries `static_assert`s on the size and offsets, compiled when `ONESHADER_CHECK_LAYOUT` is defined. - **Overlapping `packoffset` members** are an error (OSH0332) instead of invalid output. - **Push constants** follow different rules (a `float3` may share a row with the next scalar). See [Push constants on every target](https://recursiveai.org/oneshader/cookbook/push-constants/). The full rules, structured buffers and resource arrays included, are in [Generated bindings](https://recursiveai.org/oneshader/docs/generated-bindings/). Direct3D and Vulkan offsets side by side: [HLSL constant buffer packing rules](https://recursiveai.org/oneshader/articles/hlsl-constant-buffer-packing-rules/). # Push constants on Vulkan, Metal and WebGPU Declare an HLSL push-constant block once and get the C# struct plus where each API takes its bytes: a Vulkan push-constant range, a Metal `setVertexBytes` index, and a WebGPU uniform buffer, since WebGPU has no push constants. A struct-typed global or `ConstantBuffer` marked `[[vk::push_constant]]` is a push-constant block, not a descriptor. `Shaders/Mesh.hlsl`: ```hlsl struct PerDraw { float4x4 Model; float4 Tint; uint Index; }; [[vk::push_constant]] ConstantBuffer Draw; cbuffer Camera : register(b0) { float4x4 ViewProjection; }; struct VSOut { float4 Position : SV_Position; float4 Color : COLOR0; }; [shader("vertex")] VSOut VSMain(float3 position : POSITION) { VSOut o; o.Position = mul(ViewProjection, mul(Draw.Model, float4(position, 1))); o.Color = Draw.Tint; return o; } [shader("pixel")] float4 PSMain(VSOut i) : SV_Target { return i.Color; } ``` The project references `RecursiveAI.OneShader.Build` 1.2.0 with `OneShaderTargets` set to `SpirV;Msl;Wgsl`. `Program.cs`: ```csharp using System.Numerics; using System.Runtime.InteropServices; using Push.Shaders; var draw = new Mesh.PushConstants { Model = Matrix4x4.CreateTranslation(0, 1, 0), Tint = new Vector4(1, 0.5f, 0, 1), Index = 7, }; ReadOnlySpan bytes = MemoryMarshal.AsBytes(new ReadOnlySpan(in draw)); Console.WriteLine($"Push constants: {Mesh.PushConstants.Size} bytes, {bytes.Length} in the struct"); Console.WriteLine($"Vulkan: vkCmdPushConstants(offset 0, size {Mesh.Bindings.PushConstants.Size}, stages {Mesh.Bindings.PushConstants.Stages})"); Console.WriteLine($"Metal: setVertexBytes at buffer index {Mesh.Bindings.PushConstants.MslIndex}"); Console.WriteLine($"WebGPU: uniform buffer at @group({Mesh.Bindings.PushConstants.WgslGroup}) @binding({Mesh.Bindings.PushConstants.WgslBinding})"); Console.WriteLine($"Camera: set {Mesh.Bindings.Camera.Set}, binding {Mesh.Bindings.Camera.Binding}, Metal buffer {Mesh.Bindings.Camera.MslIndex}"); ``` ```text Push constants: 96 bytes, 96 in the struct Vulkan: vkCmdPushConstants(offset 0, size 96, stages Vertex) Metal: setVertexBytes at buffer index 0 WebGPU: uniform buffer at @group(0) @binding(1) Camera: set 0, binding 0, Metal buffer 1 ``` The SPIR-V of both entry points passes `spirv-val`, and the WGSL passes naga. ## How each API gets the bytes - **Vulkan:** `vkCmdPushConstants(cmd, layout, stages, 0, Size, &data)`, where `stages` is `Bindings.PushConstants.Stages` cast to `VkShaderStageFlags`, with a push-constant range at offset 0, `Size` bytes, those stages, in the pipeline layout. `Stages` is only the stages that read the block: here the vertex shader, because the pixel shader doesn’t touch `Draw`. - **Metal:** the block is a `constant T&` argument. Pass the struct’s bytes with `setVertexBytes` / `setFragmentBytes` (or `setBytes` for compute) at `MslIndex`. It takes the next buffer index in declaration order, so `Camera` moves to buffer 1. - **WebGPU:** a uniform buffer at `@group(WgslGroup) @binding(WgslBinding)`, at least `Size` rounded up to 16 bytes. Write the bytes into it before each draw. Its binding is the one after the highest binding of the highest group the shader uses, so it never collides with the shader’s own. - **One struct for all three.** The members follow the std430-style rules DXC uses for push constants (a `float3` may share a 16-byte row with the next scalar), not the constant-buffer rules, and the same bytes are valid on every target. ## Limits - Vulkan only guarantees 128 bytes of push constants. A larger block compiles with warning OSH0334. - A file may declare several blocks as long as no entry point reads more than one (OSH0333). With several, each struct is named `PushConstants`. - The block isn’t in `Bindings.All`, because it isn’t a descriptor. `Pipeline.ToVulkan()` lists it in `PushConstantRanges`, and `TryToWebGpu()` adds the stand-in uniform buffer to its bind group. The rules are in [Generated bindings](https://recursiveai.org/oneshader/docs/generated-bindings/). For ordinary constant buffers, see [Match constant buffers with C# structs](https://recursiveai.org/oneshader/cookbook/cbuffer-packing-csharp-structs/). # Shader variants with Axes in dotnet build Compile every combination of an HLSL shader’s keywords during `dotnet build` with `Axes`, or hand-picked define sets with `OneShaderVariant`, and pick a variant at run time with a typed key and no allocation. A material shader often has keywords such as alpha test on or off, and a lighting model. In `RecursiveAI.OneShader.Build` 1.2.0, one item metadata compiles the cartesian product of the keywords, once per variant and target, in parallel. The generated class gets a typed `Key`, a `Variants` table and one static class per variant. The shader is the repository’s `samples/Variants/Shaders/Material.hlsl`: ```hlsl // A sprite material with two keyword axes. The build compiles it once per combination, defining // ALPHA_TEST (a flag: defined as 1, or not defined at all) // LIGHTING (the value: NONE, LAMBERT or PBR) and LIGHTING_NONE / LIGHTING_LAMBERT / LIGHTING_PBR as 1. cbuffer Frame : register(b0) { float4x4 ViewProjection; float4 Tint; }; Texture2D Atlas : register(t1); SamplerState AtlasSampler : register(s2); #if LIGHTING_LAMBERT || LIGHTING_PBR Texture2D NormalMap : register(t3); #endif struct VSIn { float2 Position : POSITION; float2 Uv : TEXCOORD0; float4 Color : COLOR0; }; struct VSOut { float4 Position : SV_Position; float2 Uv : TEXCOORD0; float4 Color : COLOR0; }; [shader("vertex")] VSOut VSMain(VSIn i) { VSOut o; o.Position = mul(ViewProjection, float4(i.Position, 0, 1)); o.Uv = i.Uv; o.Color = i.Color; return o; } [shader("pixel")] float4 PSMain(VSOut i) : SV_Target { float4 c = Atlas.Sample(AtlasSampler, i.Uv) * i.Color * Tint; #if ALPHA_TEST clip(c.a - 0.5); #endif #if LIGHTING_LAMBERT float3 n = NormalMap.Sample(AtlasSampler, i.Uv).xyz * 2.0 - 1.0; c.rgb *= saturate(dot(n, normalize(float3(0.3, 0.5, 0.8)))); #elif LIGHTING_PBR float3 n = NormalMap.Sample(AtlasSampler, i.Uv).xyz * 2.0 - 1.0; float3 l = normalize(float3(0.3, 0.5, 0.8)); float d = saturate(dot(n, l)); c.rgb = c.rgb * d + pow(d, 16.0) * 0.04; #endif return c; } ``` A second shader, `Shaders/Sky.hlsl`, gets two named define sets instead: ```hlsl cbuffer SkyParams : register(b0) { float3 SunDirection; float Time; }; [shader("pixel")] float4 PSMain(float3 dir : TEXCOORD0) : SV_Target { float sun = saturate(dot(normalize(dir), SunDirection)); #if NIGHT float3 c = float3(0.01, 0.02, 0.05); #ifdef STARS c += step(0.999, frac(sin(dot(dir.xy, float2(12.9898, 78.233))) * 43758.5453)); #endif #else float3 c = lerp(float3(0.4, 0.6, 0.9), float3(1.0, 0.9, 0.7), sun); #endif return float4(c, 1); } ``` The project: ```xml Exe net10.0 enable enable ``` `Program.cs`: ```csharp using Axes.Shaders; Console.WriteLine($"{Material.Variants.Count} variants"); foreach (Material.Variant v in Material.Variants.All) Console.WriteLine($"{v.Index} {v.Name,-28} {v.Defines,-29} pixel {v.PSMain.SpirV.Length} bytes, {v.Bindings.Length} bindings"); // A typed key picks one variant: an array index, no allocation. var key = new Material.Key(alphaTest: true, lighting: Material.LightingValue.Pbr); Material.Variant pbr = Material.Variants.Get(key); Console.WriteLine($"Get({key.AlphaTest}, {key.Lighting}) -> {pbr.Name}"); foreach (var b in pbr.Bindings) Console.WriteLine($" {b.Name}: set {b.Set}, binding {b.Binding}"); // Every variant has the same Frame layout, so the struct is generated once. var frame = new Material.Frame { Tint = System.Numerics.Vector4.One }; Console.WriteLine($"Frame: {Material.Frame.Size} bytes, shared by all variants"); // The same variant as a static class, shaped like a file without variants. Console.WriteLine($"Material.Variants.AlphaTestOn_LightingPbr.PSMain: {Material.Variants.AlphaTestOn_LightingPbr.PSMain.SpirV.Length} bytes"); // Named variants: one per OneShaderVariant item. foreach (Sky.Variant v in Sky.Variants.All) Console.WriteLine($"Sky {v.Name}: {v.Defines}, pixel {v.PSMain.SpirV.Length} bytes"); Sky.Variant night = Sky.Variants.Get(new Sky.Key(Sky.VariantName.Night)); Console.WriteLine($"Get(Night) -> {night.Name}"); ``` ```text 6 variants 0 AlphaTestOff_LightingNone LIGHTING=NONE pixel 1248 bytes, 3 bindings 1 AlphaTestOff_LightingLambert LIGHTING=LAMBERT pixel 1956 bytes, 4 bindings 2 AlphaTestOff_LightingPbr LIGHTING=PBR pixel 2080 bytes, 4 bindings 3 AlphaTestOn_LightingNone ALPHA_TEST=1 LIGHTING=NONE pixel 1396 bytes, 3 bindings 4 AlphaTestOn_LightingLambert ALPHA_TEST=1 LIGHTING=LAMBERT pixel 2072 bytes, 4 bindings 5 AlphaTestOn_LightingPbr ALPHA_TEST=1 LIGHTING=PBR pixel 2196 bytes, 4 bindings Get(True, Pbr) -> AlphaTestOn_LightingPbr Frame: set 0, binding 0 Atlas: set 0, binding 1 AtlasSampler: set 0, binding 2 NormalMap: set 0, binding 3 Frame: 80 bytes, shared by all variants Material.Variants.AlphaTestOn_LightingPbr.PSMain: 2196 bytes Sky Day: NIGHT=0, pixel 1092 bytes Sky Night: NIGHT=1 STARS, pixel 864 bytes Get(Night) -> Night ``` ## How the axes become defines - **A flag** (`ALPHA_TEST`) is either defined as `1` or not defined at all, so test it with `#if` or `#ifdef`. - **A choice** (`LIGHTING=NONE|LAMBERT|PBR`) defines `LIGHTING` as the value, and `LIGHTING_PBR` (for example) as `1`, so `#if LIGHTING_PBR` works. - **Order** is the cartesian product with the first axis changing slowest. `Key.Index` is the position in `Variants.All`. - **Names** are `On`/`Off` for a flag and `` for a choice, PascalCased and joined by `_`. Choices become enums (`Material.LightingValue`), flags `bool`s, and named sets `VariantName`. - **Bindings can differ.** The normal map exists only in the lit variants, so read slots from the variant (`v.Bindings`, `v.Pipeline`), not from the file. A struct that every variant lays out the same way (`Frame` here) is generated once. If a member sits under `#if`, each variant gets its own struct instead. - **Other metadata** (`Defines`, `Targets`, `$(OneShaderDefines)`) applies to every variant. Use either `Axes` or `OneShaderVariant` items on one shader, not both. ## Building incrementally Only the variants whose define set, shader, includes or options changed are compiled again. With `dotnet build -v:d`, a build with nothing changed logs: ```text OneShader: 0 compiled, 8 up to date. ``` After adding `TOON` to the lighting axis (`LIGHTING=NONE|LAMBERT|PBR|TOON`), just the two new combinations compile: ```text OneShader: 2 compiled, 8 up to date. ``` ## Errors under some variants An error is reported once, on the HLSL line, with the variants it happened in. Misspelling `d` in the PBR branch gives: ```text /tmp/oneshader-docs/cookbook/Axes/Shaders/Material.hlsl(38,29): error OSH0300: 'dd' is not declared. Did you mean 'd'? (variants LIGHTING=PBR; ALPHA_TEST=1, LIGHTING=PBR) [/tmp/oneshader-docs/cookbook/Axes/Axes.csproj] ``` ## Limits - **64 variants per shader** by default. More is error OSH0905 naming the fix: raise it with `MaxVariants="256"` on the item or `` in the project. Each variant is compiled for every selected target, so the product grows fast. - **A cbuffer, struct or entry point can’t share the file’s name.** A `cbuffer Sky` in `Sky.hlsl` is error OSH0951, as in any file. - For keyword values that change per draw without recompiling, use [specialization constants](https://recursiveai.org/oneshader/cookbook/specialization-constants/) instead. To compile variants at run time, see [Shader variants with defines and includes](https://recursiveai.org/oneshader/cookbook/shader-variants-defines-includes/) and [Cache shader compiles](https://recursiveai.org/oneshader/cookbook/shader-compile-cache/). The generated members are listed under Variants in [Generated bindings](https://recursiveai.org/oneshader/docs/generated-bindings/), and the MSBuild side in [the build package](https://recursiveai.org/oneshader/docs/build-package/). # Specialization constants on Vulkan, Metal and WebGPU Declare HLSL specialization constants once with `[[vk::constant_id]]` and set them when the pipeline is created: `VkSpecializationInfo` on Vulkan, function constants on Metal and pipeline-overridable constants on WebGPU, from one compiled shader. A specialization constant is a value the shader reads like a constant but the pipeline chooses, so one compiled shader serves many variants without recompiling. OneShader 1.2.0 compiles `[[vk::constant_id(N)]] const T Name = default;` (`bool`, `int`, `uint` or `float`) to each target’s mechanism, and generates the C# to fill them in. `Shaders/Blur.hlsl`, a separable Gaussian blur whose radius, direction and sigma are pipeline-time constants: ```hlsl // Pipeline-time constants: one compiled shader, many pipelines. [[vk::constant_id(0)]] const int Radius = 4; [[vk::constant_id(1)]] const bool Horizontal = true; [[vk::constant_id(2)]] const float Sigma = 2.0; Texture2D Source : register(t0); [[vk::image_format("rgba8")]] RWTexture2D Target : register(u1); [shader("compute")] [numthreads(8, 8, 1)] void CSMain(uint3 id : SV_DispatchThreadID) { int2 step = Horizontal ? int2(1, 0) : int2(0, 1); float4 sum = 0; float total = 0; for (int i = -Radius; i <= Radius; i++) { float w = exp(-(i * i) / (2 * Sigma * Sigma)); sum += Source[int2(id.xy) + step * i] * w; total += w; } Target[id.xy] = sum / total; } ``` The project references `RecursiveAI.OneShader.Build` 1.2.0 with `OneShaderTargets` set to `SpirV;Msl;Wgsl`. `Program.cs`: ```csharp using OneShader.Layouts; using OneShader.Layouts.Metal; using OneShader.Layouts.Vulkan; using Spec.Shaders; // The constants, their ids and defaults, generated from the shader. Console.WriteLine($"Radius: id {Blur.SpecializationConstants.Radius.Id}, default {Blur.SpecializationConstants.Radius.Default}"); Console.WriteLine($"CSMain reads ids [{string.Join(", ", Blur.CSMain.SpecConstantIds.ToArray())}]"); // A vertical blur with radius 8: set fields on the generated struct. Blur.Specialization s = Blur.Specialization.Defaults; s.Radius = 8; s.Horizontal = 0; // a bool is a 4-byte 0 or 1 // Vulkan: the contents of VkSpecializationInfo. SpecializationInfo vk = s.ToVulkan(Blur.CSMain.SpecConstantIds); foreach (SpecializationMapEntry e in vk.MapEntries) Console.WriteLine($"Vulkan: constantID {e.ConstantId}, offset {e.Offset}, size {e.Size}"); Console.WriteLine($"Vulkan: data {Convert.ToHexString(vk.Data)}"); // Metal: MTLFunctionConstantValues entries, by index and MTLDataType. foreach (FunctionConstantValue v in s.ToMetal(Blur.CSMain.SpecConstantIds)) Console.WriteLine($"Metal: index {v.Index}, MTLDataType {v.Type.DataType}, bits 0x{v.Bits:x8}"); // WebGPU: the compute stage's constants record, keyed by decimal id. Console.WriteLine($"WebGPU: {Record(s.ToWebGpu(Blur.CSMain.SpecConstantIds))}"); // The same by name, from the entry layout (this also works for shaders compiled at run time). SpecializationValues byName = SpecializationValues.For(Blur.CSMain.Layout).Set("Radius", 8).Set("Horizontal", false); Console.WriteLine($"By name: {Record(byName.ToWebGpu())}"); // What each target compiled the constant to. Console.WriteLine(Line(Blur.CSMain.Wgsl, "override Radius")); Console.WriteLine(Line(Blur.CSMain.Msl, "function_constant(0)")); static string Record(IReadOnlyDictionary c) => "{ " + string.Join(", ", c.Select(kv => $"\"{kv.Key}\": {kv.Value}")) + " }"; static string Line(ReadOnlySpan utf8, string text) => System.Text.Encoding.UTF8.GetString(utf8).Split('\n').First(l => l.Contains(text)).Trim(); ``` ```text Radius: id 0, default 4 CSMain reads ids [0, 1, 2] Vulkan: constantID 0, offset 0, size 4 Vulkan: constantID 1, offset 4, size 4 Vulkan: constantID 2, offset 8, size 4 Vulkan: data 080000000000000000000040 Metal: index 0, MTLDataType 29, bits 0x00000008 Metal: index 1, MTLDataType 53, bits 0x00000000 Metal: index 2, MTLDataType 3, bits 0x40000000 WebGPU: { "0": 8, "1": 0, "2": 2 } By name: { "0": 8, "1": 0, "2": 2 } @id(0) override Radius: i32 = 4i; constant int Radius [[function_constant(0)]]; ``` The SPIR-V passes `spirv-val` and the WGSL passes naga. On SPIR-V the constants are `OpSpecConstant` (and `OpSpecConstantTrue` for the `bool`) with `SpecId` 0, 1 and 2, the ids DXC writes for the same source. ## Passing them to each API - **Vulkan:** `MapEntries` and `Data` are the contents of `VkSpecializationInfo` (`pMapEntries`, `pData`). Each constant is 4 bytes, in declaration order: `08000000` is the radius 8, `00000000` the `false`, `00000040` the float 2.0. - **Metal:** set each value on an `MTLFunctionConstantValues` at its index, with the `MTLDataType` number (29 is `int`, 53 `bool`, 3 `float`). A function that declares function constants has to be created with `makeFunction(name:constantValues:)`, even to get the defaults: pass an empty `MTLFunctionConstantValues` for those. The MSL reads each constant through `is_function_constant_defined`, with the default as the fallback. - **WebGPU:** the record goes in the stage’s `constants` when creating the pipeline. Pass the entry point’s `SpecConstantIds`, as above, because WebGPU rejects constants the entry point doesn’t declare. Vulkan and Metal ignore extra ones. - **By name:** `SpecializationValues.For(layout)` takes an `EntryLayout` from the generated code or from run-time reflection (`EntryLayout.TryCreate`), with typed `Set` overloads. A value of the wrong type or an unknown name throws. ## Fixing constants at compile time To bake some values in instead, compile with `CompileOptions.SpecializationOverrides` (`RecursiveAI.OneShader` 1.2.0): ```csharp using OneShader; // Fix Radius at compile time: it becomes a plain constant that the optimizer folds. // The others stay pipeline-time constants. CompileResult result = ShaderCompiler.CompileFile("Shaders/Blur.hlsl", new CompileOptions { SpecializationOverrides = [new("Radius", 2), new("Blurr", 1)], }); foreach (var d in result.Diagnostics) Console.WriteLine(d); Console.WriteLine($"left as constants: {string.Join(", ", result.Reflection!.SpecConstants.Select(c => $"{c.Name} (id {c.Id})"))}"); CompileResult bad = ShaderCompiler.CompileFile("Shaders/Blur.hlsl", new CompileOptions { SpecializationOverrides = [new("Radius", 2.5)], }); foreach (var d in bad.Diagnostics) Console.WriteLine(d); ``` ```text Shaders/Blur.hlsl(1,1): warning OSH0335: The shader declares no specialization constant named 'Blurr', so the override is ignored. Check the spelling against the [[vk::constant_id]] declarations. left as constants: Horizontal (id 1), Sigma (id 2) Shaders/Blur.hlsl(1,1): error OSH0335: The value 2.5 does not fit the int specialization constant 'Radius': use a whole number in its range. ``` ## Limits - **Types:** `bool`, `int`, `uint` and `float` scalars. A vector or struct, `static`, a missing or non-constant default, a negative id or a duplicate id is error OSH0335. - **Not where the size must be known at compile time:** an array size, a `[numthreads]` argument or a `case` label that uses one is error OSH0310, as DXC refuses the same uses: ```text /tmp/nt.hlsl(3,33): error OSH0310: Argument 1 of 'numthreads' must be an integer constant. It uses a specialization constant, whose value is only known when the pipeline is created: use a literal or a 'static const' there, or choose the value with a #define. ``` - **Hot reload:** `InterfaceCompatibility` reports a removed constant or a changed id or type. New constants and changed defaults stay compatible. For variants that change resources or code structure, compile them separately with [Axes](https://recursiveai.org/oneshader/cookbook/shader-variants-axes/). The generated members are under Specialization constants in [Generated bindings](https://recursiveai.org/oneshader/docs/generated-bindings/). # 16-bit and 64-bit types: half, int16_t, double and uint64_t Use native 16-bit `half` and `int16_t` and 64-bit `double` and `uint64_t` in HLSL with OneShader 1.2.0, get the matching C# structs, and see which device features each target needs and what Metal and WebGPU refuse. OneShader 1.2.0 compiles the sized types the way DXC does. With 16-bit types enabled (DXC’s `-enable-16bit-types`), `half`, `min16float`, `min16int`, `min16uint` and the `float16_t`, `int16_t` and `uint16_t` families are native 16-bit types, in arithmetic and in buffers with DXC’s byte layout. `double`, `int64_t` and `uint64_t` work on SPIR-V in arithmetic, in buffers and in 64-bit `Interlocked*` operations. ## The project 16-bit types are switched on for the project. `Totals.hlsl` uses 64-bit types, which only SPIR-V has, so that file is built for SPIR-V only: ```xml Exe net10.0 enable enable true SpirV;Msl;Wgsl ``` `Shaders/Particles.hlsl` keeps particle state in half-precision floats: ```hlsl struct Particle { half3 Velocity; // 16-bit floats: half the bytes of a float3 half Life; uint Flags; }; cbuffer Step : register(b0) { half DeltaTime; uint Count; }; RWStructuredBuffer Particles : register(u1); [shader("compute")] [numthreads(64, 1, 1)] void CSMain(uint3 id : SV_DispatchThreadID) { if (id.x >= Count) return; Particle p = Particles[id.x]; p.Life -= DeltaTime; p.Velocity *= (half)0.98; if (p.Life <= 0) p.Flags |= 1; Particles[id.x] = p; } ``` `Shaders/Totals.hlsl` keeps a 64-bit byte counter and a double-precision sum: ```hlsl struct Sums { uint64_t Bytes; // a 64-bit counter that does not wrap at 4 GB double Energy; // 64-bit float accumulation }; StructuredBuffer Sizes : register(t0); StructuredBuffer Samples : register(t1); RWStructuredBuffer Result : register(u2); [shader("compute")] [numthreads(64, 1, 1)] void CSMain(uint3 id : SV_DispatchThreadID) { InterlockedAdd(Result[0].Bytes, (uint64_t)Sizes[id.x]); double e = (double)Samples[id.x] * (double)Samples[id.x]; if (id.x == 0) Result[0].Energy += e; } ``` `Program.cs`: ```csharp using System.Runtime.InteropServices; using OneShader.Layouts.WebGpu; using Sized.Shaders; // 16-bit members are System.Half, one field per component. var p = new Particles.Particle { Velocity_X = (Half)1.5f, Velocity_Y = (Half)0, Velocity_Z = (Half)(-2), Life = (Half)3 }; Console.WriteLine($"Particle: {Particles.Particle.Size} bytes, Life at {Marshal.OffsetOf("Life")}, Flags at {Marshal.OffsetOf("Flags")}"); Console.WriteLine($"Step: {Particles.Step.Size} bytes, Count at {Marshal.OffsetOf("Count")}"); // 64-bit members are ulong and double. Console.WriteLine($"Sums: {Totals.Sums.Size} bytes, Energy at {Marshal.OffsetOf("Energy")}"); // What the device must enable. Console.WriteLine($"Particles.CSMain.RequiredFeatures: {Particles.CSMain.RequiredFeatures}"); Console.WriteLine($"Totals.CSMain.RequiredFeatures: {Totals.CSMain.RequiredFeatures}"); if (Particles.Pipeline.TryToWebGpu(out WebGpuPipelineLayout? webgpu, out var error)) Console.WriteLine($"WebGPU requiredFeatures: [{string.Join(", ", webgpu!.RequiredFeatures ?? [])}]"); else Console.WriteLine(error); // half in WGSL is f16. string wgsl = System.Text.Encoding.UTF8.GetString(Particles.CSMain.Wgsl); Console.WriteLine(wgsl.Split('\n')[0]); ``` ```text Particle: 16 bytes, Life at 6, Flags at 8 Step: 16 bytes, Count at 4 Sums: 16 bytes, Energy at 8 Particles.CSMain.RequiredFeatures: Float16, StorageBuffer16, UniformBuffer16 Totals.CSMain.RequiredFeatures: Float64, Int64, Int64Atomics WebGPU requiredFeatures: [shader-f16] enable f16; ``` The layouts are DXC’s. `dxc -enable-16bit-types -spirv` on the same `Particles.hlsl` gives member offsets 0, 6 and 8 and an array stride of 16, and `dxc -spirv` on `Totals.hlsl` gives offsets 0 and 8 and a stride of 16, as OneShader does. Both SPIR-V files pass `spirv-val`, and the WGSL passes naga. ## What the generated C# looks like - `half` is `System.Half`, `int16_t` `short`, `uint16_t` `ushort`, `int64_t` `long`, `uint64_t` `ulong`, `double` `double`. - Vectors of these have no `System.Numerics` type, so they’re one field per component: `half3 Velocity` becomes `Velocity_X`, `Velocity_Y` and `Velocity_Z`. A matrix of a 16-bit or 64-bit element is an error that says to split it. - **`RequiredFeatures`** is on every entry-point class and its `Layout`. Enable those device features when creating the Vulkan device (`PipelineLayoutInfo.RequiredFeatures` lists them too), and request `shader-f16` on WebGPU. - Turn 16-bit types on per project with ``, per file with `Enable16BitTypes="true"` metadata, with `-enable-16bit-types` on the command line, or with `CompileOptions.Enable16BitTypes` at run time. ## Without 16-bit types As in DXC, `half` and the `min16*` types are then 32-bit values (marked `RelaxedPrecision` on SPIR-V), and the explicitly sized names are an error: ```text i16.hlsl(1,20): error OSH0360: The type 'uint16_t' needs 16-bit types enabled. Compile with -enable-16bit-types (MSBuild: OneShaderEnable16BitTypes=true), or use half, min16float, int or uint. ``` A `double` passed to a float-only intrinsic is narrowed to `float`, as DXC does, with a warning: ```text dn.hlsl(2,64): warning OSH0361: 'sqrt' has no double version: the argument is converted to float, so the result has float precision. Use +, -, *, /, min, max, abs, clamp or fma to keep double precision. ``` ## What each target has | Type | SPIR-V | Metal | WGSL | |---|---|---|---| | `half` | yes | yes (native `half`) | yes (`f16`, `enable f16;`) | | `int16_t`, `uint16_t` | yes | yes (`short`, `ushort`) | no | | `int64_t`, `uint64_t` | yes | yes (`long`, `ulong`), not as vectors in buffers | no | | `double` | yes | no | no | | 64-bit `Interlocked*` | yes, on structured buffers and `groupshared` | only unsigned `InterlockedMax`/`Min` on a buffer, result unused | no | A refusal is an error at the first use, only when that target is selected, with what to change. `Totals.hlsl` compiled for Metal and WGSL: ```text /tmp/oneshader-docs/cookbook/Sized/Shaders/Totals.hlsl(15,5): error OSH5102: MSL has no 64-bit floating-point type. Use float, or compile this shader for SPIR-V only. /tmp/oneshader-docs/cookbook/Sized/Shaders/Totals.hlsl(15,5): error OSH5102: Metal 2.4 supports 64-bit atomics only as unsigned InterlockedMax / InterlockedMin on a buffer without using the returned value. Use that form, 32-bit atomics, or compile this shader for SPIR-V only. /tmp/oneshader-docs/cookbook/Sized/Shaders/Totals.hlsl(15,5): error OSH5202: WGSL has no 64-bit floating-point type. Use float, or compile this shader for SPIR-V only. /tmp/oneshader-docs/cookbook/Sized/Shaders/Totals.hlsl(15,5): error OSH5202: WGSL has no 64-bit integer types. Use int or uint (or a uint2 pair), or compile this shader for SPIR-V or MSL only. ``` And a `uint16_t` buffer on WGSL: ```text /tmp/i16.hlsl(2,57): error OSH5202: WGSL has no 16-bit integer types (only f16). Use int or uint, or compile this shader for SPIR-V or MSL only. ``` The per-target rules are in [SPIR-V](https://recursiveai.org/oneshader/docs/spirv-backend/), [MSL](https://recursiveai.org/oneshader/docs/msl-backend/) and [WGSL](https://recursiveai.org/oneshader/docs/wgsl-backend/), and the generated members under 16-bit and 64-bit members in [Generated bindings](https://recursiveai.org/oneshader/docs/generated-bindings/). # Read shader reflection in C#: bindings, buffer layouts and vertex inputs Read what a compiled HLSL shader binds, where, in which stages, its buffer layouts and its vertex inputs, from C# or as JSON, and turn it into Vulkan descriptor set layouts, with no native SPIRV-Reflect. Every compile returns reflection next to the code: resources with their slots per target and their byte layouts, push-constant blocks, and each entry point’s inputs, outputs, used resources and thread-group size. It comes from the compiler, not from parsing the SPIR-V afterwards, and it’s the same model for all three targets. ## From C# With `RecursiveAI.OneShader` 1.2.0 and the `Shaders/Sprite.hlsl` of [Getting started](https://recursiveai.org/oneshader/docs/getting-started/): ```csharp using OneShader; using OneShader.Layouts; using OneShader.Layouts.Vulkan; using OneShader.Reflection; CompileResult result = ShaderCompiler.CompileFile("Shaders/Sprite.hlsl"); if (!result.Success) { foreach (var d in result.Diagnostics) Console.WriteLine(d); return; } ShaderReflection reflection = result.Reflection!; // Resources: kind, slot, stages, and the byte layout of buffers. foreach (ResourceInfo r in reflection.Resources) { TargetSlot slot = r.Slots.First(s => s.Target == CompileTarget.SpirV); Console.WriteLine($"{r.Name}: {r.Kind}, set {slot.Set} binding {slot.Binding}, stages {r.Stages}"); if (r.Layout is { Kind: ReflectedTypeKind.Struct } layout) foreach (ReflectedMember m in layout.Members) Console.WriteLine($" {m.Name} at {m.Offset}, {m.Type.Size} bytes"); } // Entry points: stage, resources used, and vertex inputs with their formats. foreach (EntryPointInfo e in reflection.EntryPoints) { Console.WriteLine($"{e.Name}: {e.Stage}, uses {string.Join(", ", e.UsedResources)}"); foreach (StageVariable v in e.Inputs.Where(v => v.Format is not null)) Console.WriteLine($" in location {v.Location}: {v.Semantic}{v.SemanticIndex} {v.Format}"); } // From reflection to a Vulkan pipeline layout. EntryLayout.TryCreate(reflection, "VSMain", out EntryLayout? vs, out _); EntryLayout.TryCreate(reflection, "PSMain", out EntryLayout? ps, out _); if (!PipelineLayout.TryCreate([vs!, ps!], out PipelineLayout? pipeline, out LayoutError? error)) { Console.WriteLine(error); return; } foreach (DescriptorSetLayoutInfo set in pipeline!.ToVulkan().SetLayouts) foreach (DescriptorSetLayoutBindingInfo b in set.Bindings) Console.WriteLine($"VkDescriptorSetLayoutBinding {{ binding = {b.Binding}, descriptorType = {(uint)b.DescriptorType}, descriptorCount = {b.DescriptorCount}, stageFlags = 0x{(uint)b.StageFlags:x} }}"); ``` ```text Frame: ConstantBuffer, set 0 binding 0, stages Vertex ViewProjection at 0, 64 bytes Tint at 64, 16 bytes Atlas: SampledTexture, set 0 binding 1, stages Pixel AtlasSampler: Sampler, set 0 binding 2, stages Pixel VSMain: Vertex, uses Frame in location 0: POSITION0 Float32x2 in location 1: TEXCOORD0 Float32x2 in location 2: COLOR0 Unorm8x4 PSMain: Pixel, uses Atlas, AtlasSampler VkDescriptorSetLayoutBinding { binding = 0, descriptorType = 6, descriptorCount = 1, stageFlags = 0x1 } VkDescriptorSetLayoutBinding { binding = 1, descriptorType = 2, descriptorCount = 1, stageFlags = 0x10 } VkDescriptorSetLayoutBinding { binding = 2, descriptorType = 0, descriptorCount = 1, stageFlags = 0x10 } ``` `descriptorType` 6, 2 and 0 are `VK_DESCRIPTOR_TYPE_UNIFORM_BUFFER`, `SAMPLED_IMAGE` and `SAMPLER`. `stageFlags` 0x1 and 0x10 are `VK_SHADER_STAGE_VERTEX_BIT` and `FRAGMENT_BIT`. - `Slots` has one entry per compiled target. Vulkan and WGSL use the same set and binding, and the Metal slot is the argument index. - `Layout` is the byte layout of a buffer, members, offsets, array and matrix strides included, equal for every target. - `PipelineLayout.TryCreate` merges entry points that run together: a binding used by both gets the union of their stages, and a conflict comes back as a `LayoutError` value. Unused sets are empty, as Vulkan needs them. - Projects that use `RecursiveAI.OneShader.Build` don’t need this at run time: the generated class has the same data as constants, and `Sprite.Reflection` holds the reflection it was generated from. ## As JSON from the command line ```sh oneshader -T vs_6_0 -E VSMain --reflect sprite.vs.json Shaders/Sprite.hlsl -Fo sprite.vs.spv ``` The start of the entry point in `sprite.vs.json`: ```json "entryPoints": [ { "name": "VSMain", "stage": "Vertex", "vertexStride": 20, "inputs": [ { "name": "i.Position", "semantic": "POSITION", "semanticIndex": 0, "location": 0, "builtIn": "None", "scalar": "Float", "components": 2, "format": "Float32x2", "offset": 0, "interpolation": "Default" }, ``` The JSON is schema version 1: keys in camelCase, enums by name, a fixed order, `\n` line ends and no machine paths, so the same compile gives the same bytes. Breaking changes raise `schemaVersion`. The build task’s own reflection file has a different, line-based format, described in [Reflection file format](https://recursiveai.org/oneshader/docs/reflection-format/). To check a recompiled shader against the reflection the C# was built from, see [Reload shaders while the app runs](https://recursiveai.org/oneshader/cookbook/hot-reload-shaders/). # Reload shaders while the app runs Recompile an HLSL shader when its file changes, swap in the new code while the app keeps running, and find out when an edit changed the interface the generated C# was built against, such as a constant-buffer offset. A project built with `RecursiveAI.OneShader.Build` embeds its shaders and generates C# against their layout. During development you can recompile at run time with `ShaderCompiler.CompileFile`. An edit to a function body can be swapped in. An edit that moves a constant-buffer member can’t, because the generated `Frame` struct would write the wrong bytes. `InterfaceCompatibility.Compare` tells the two apart. The project references `RecursiveAI.OneShader.Build` 1.2.0 and contains the `Shaders/Sprite.hlsl` of [Getting started](https://recursiveai.org/oneshader/docs/getting-started/). `Program.cs` runs four edits through the same reload function: ```csharp using OneShader; using OneShader.Reflection; using HotReload.Shaders; // What the C# was built against: the reflection embedded by the build. ShaderReflection built = Sprite.Reflection; // 1. Recompile the shader as it is on disk (in a real loop: when the file watcher fires). Reload("Shaders/Sprite.hlsl"); // 2. An edit to a function body: still compatible. string source = File.ReadAllText("Shaders/Sprite.hlsl"); string edited = Path.Combine(Path.GetTempPath(), "Sprite.hlsl"); // outside the project, or the build compiles it too File.WriteAllText(edited, source.Replace("* i.Color;", "* i.Color * 0.5;")); Reload(edited); // 3. A new constant-buffer member before Tint: Tint moves, so the generated Frame struct is wrong. File.WriteAllText(edited, source.Replace("float4x4 ViewProjection;", "float4x4 ViewProjection; float4 Fog;")); Reload(edited); // 4. A typo: diagnostics, not an exception. File.WriteAllText(edited, source.Replace("AtlasSampler, i.Uv", "AtlasSampler, i.UV")); Reload(edited); void Reload(string path) { CompileResult result = ShaderCompiler.CompileFile(path); Console.WriteLine($"{path}: watch {string.Join(", ", result.Files.Select(Path.GetFileName))}"); if (!result.Success) { foreach (var d in result.Diagnostics) Console.WriteLine(d); // keep running the last good shader return; } CompatibilityReport report = InterfaceCompatibility.Compare(built, result.Reflection!); if (report.IsCompatible) Console.WriteLine($" compatible: swap in {result.Find("PSMain", CompileTarget.SpirV)!.Data.Length} bytes of SPIR-V"); else Console.WriteLine($" not compatible:\n{report}"); } ``` ```text Shaders/Sprite.hlsl: watch Sprite.hlsl compatible: swap in 856 bytes of SPIR-V /tmp/Sprite.hlsl: watch Sprite.hlsl compatible: swap in 920 bytes of SPIR-V /tmp/Sprite.hlsl: watch Sprite.hlsl not compatible: incompatible with the compiled-in C# bindings: Frame.Tint moved from offset 64 to 80 Frame.Fog was added at offset 64 Rebuild the project to update the generated C# bindings. /tmp/Sprite.hlsl: watch Sprite.hlsl /tmp/Sprite.hlsl(24,74): error OSH0308: The struct 'VSOut' has no member 'UV'. Did you mean 'Uv'? ``` ## Good to know - **What to watch.** `result.Files` lists the file and every `#include` it read, even when the compile fails, so the watcher keeps following a shader that’s broken mid-edit. - **Compatible** means everything the generated C# relies on is unchanged: constant- and structured-buffer layouts (offsets, types, strides, sizes), vertex input locations, formats and offsets, each resource’s set, binding, kind, count and stages, the entry points and their stages, and compute thread-group sizes. Edits to function bodies, new entry points, and resources that become visible to fewer stages are compatible. - **Errors never throw.** A failed compile returns diagnostics, so the loop keeps the last good shader running. - **Threads.** `ShaderCompiler.Compile` can be called from several threads at once, and gives the same bytes as a serial compile. - **Keep scratch copies out of the project.** Every `*.hlsl` under the project folder is compiled by the build, edited copies included. - **`Sprite.Reflection`** is stored as text in the assembly (about 1 to 4 KB per shader) and parsed on first read. A program that never reads it carries no parser after trimming. The repository’s `samples/HotReload` is a complete console loop with a file watcher. For reading the reflection itself, see [Read shader reflection in C#](https://recursiveai.org/oneshader/cookbook/spirv-reflection-csharp/). # Shader variants with defines and includes Compile several variants of one HLSL source at run time with preprocessor defines, and resolve `#include` from memory instead of the file system, with an asset pack, a database or an editor buffer as the source. `CompileOptions.Defines` is the C# form of `-D`, and `CompileOptions.IncludeResolver` is a hook the compiler calls for every `#include`. Without a resolver, `ShaderCompiler.Compile` treats `#include` as an error. `CompileFile` reads includes from disk, searching the including file’s folder and then `IncludeDirectories`. With `RecursiveAI.OneShader` 1.2.0: ```csharp using System.Security.Cryptography; using OneShader; // Shader sources held in memory: an asset pack, a database, an editor buffer. var files = new Dictionary { ["blur.hlsl"] = """ #include "common/sampling.hlsli" Texture2D Source : register(t0); SamplerState Linear : register(s1); [shader("pixel")] float4 PSMain(float2 uv : TEXCOORD0) : SV_Target { float4 sum = 0; [unroll] for (int i = -RADIUS; i <= RADIUS; i++) sum += Source.Sample(Linear, uv + TexelOffset(i)); return sum / (2 * RADIUS + 1); } """, ["common/sampling.hlsli"] = """ #ifndef HORIZONTAL #define HORIZONTAL 1 #endif float2 TexelOffset(int i) { return HORIZONTAL ? float2(i / 1920.0, 0) : float2(0, i / 1080.0); } """, }; ResolvedInclude? Resolve(string requested, string includingFile, bool isSystemInclude) => files.TryGetValue(requested, out string? text) ? new ResolvedInclude(requested, text) : null; // One source, four variants. foreach (string radius in new[] { "2", "6" }) foreach (string horizontal in new[] { "1", "0" }) { CompileResult result = ShaderCompiler.Compile(files["blur.hlsl"], "blur.hlsl", new CompileOptions { Defines = [new("RADIUS", radius), new("HORIZONTAL", horizontal)], IncludeResolver = Resolve, }); string status = result.Success ? "SHA-256 " + Convert.ToHexStringLower(SHA256.HashData(result.Find("PSMain", CompileTarget.SpirV)!.Data))[..12] : string.Join("; ", result.Diagnostics); Console.WriteLine($"RADIUS={radius} HORIZONTAL={horizontal}: {status}, read {string.Join(", ", result.Files)}"); } // A missing include is a diagnostic on the #include line. files["blur.hlsl"] = files["blur.hlsl"].Replace("common/sampling.hlsli", "common/sample.hlsli"); CompileResult missing = ShaderCompiler.Compile(files["blur.hlsl"], "blur.hlsl", new CompileOptions { Defines = [new("RADIUS", "2")], IncludeResolver = Resolve, }); Console.WriteLine(missing.Diagnostics[0]); ``` ```text RADIUS=2 HORIZONTAL=1: SHA-256 605dd910624c, read blur.hlsl, common/sampling.hlsli RADIUS=2 HORIZONTAL=0: SHA-256 2fcade900707, read blur.hlsl, common/sampling.hlsli RADIUS=6 HORIZONTAL=1: SHA-256 0479e8957f19, read blur.hlsl, common/sampling.hlsli RADIUS=6 HORIZONTAL=0: SHA-256 bd1d621c19f2, read blur.hlsl, common/sampling.hlsli blur.hlsl(1,10): error OSH0120: cannot find include file 'common/sample.hlsli'; searched: 'the include resolver'. Check the spelling, or add the directory that contains it to the include paths ``` ## Good to know - **Defines.** Each entry is a name and a value. A `null` value defines the name as 1, like `-D NAME`. Defines apply before the first line of the source, so `#ifndef` defaults in an include work as above. - **The resolver** gets the requested name, the including file and whether it was `<...>` (system) or `"..."`. Return a `ResolvedInclude(path, text)`, or `null` for “not found”. The path you return is what diagnostics and `result.Files` show, and it’s the `includingFile` the resolver gets for includes nested inside that file, so relative includes can be resolved from it. - **Variants are deterministic.** The same source, defines and options always give the same bytes, so a hash of the inputs is a safe cache key. - **Errors in an include** point at the include’s own line, under the path the resolver returned. - **The same options** exist on the command line (`-D`, `-I`) and in the build package (`Defines` and `IncludeDirectories` metadata, or `OneShaderDefines` and `OneShaderIncludeDirectories`). See [Compile HLSL in a C# project](https://recursiveai.org/oneshader/cookbook/compile-hlsl-csharp/). # Cache shader compiles at run time Keep compiled HLSL variants in memory with `ShaderCompileCache`, so a material system or editor that asks for the same shader again gets it without recompiling. An edited include is noticed, and concurrent requests for one variant compile it once. `ShaderCompileCache`, new in `RecursiveAI.OneShader` 1.2.0, answers a compile from memory when the same inputs were compiled before. The key is the source text, the path, every option (targets, defines, include directories, entry points, matrix packing, 16-bit types, optimization, names) and the content of every include the earlier compile read. Includes are looked up again on each call, so editing one makes the next call compile. The shader is the material from [Shader variants with Axes](https://recursiveai.org/oneshader/cookbook/shader-variants-axes/), with the light direction moved into an include. `Shaders/Common/Lighting.hlsli`: ```hlsl static const float3 LightDir = normalize(float3(0.3, 0.5, 0.8)); ``` `Shaders/Material.hlsl`, from its fourth line: ```hlsl #include "Lighting.hlsli" cbuffer Frame : register(b0) { float4x4 ViewProjection; float4 Tint; }; Texture2D Atlas : register(t1); SamplerState AtlasSampler : register(s2); #if LIGHTING_LAMBERT || LIGHTING_PBR Texture2D NormalMap : register(t3); #endif struct VSIn { float2 Position : POSITION; float2 Uv : TEXCOORD0; float4 Color : COLOR0; }; struct VSOut { float4 Position : SV_Position; float2 Uv : TEXCOORD0; float4 Color : COLOR0; }; [shader("vertex")] VSOut VSMain(VSIn i) { VSOut o; o.Position = mul(ViewProjection, float4(i.Position, 0, 1)); o.Uv = i.Uv; o.Color = i.Color; return o; } [shader("pixel")] float4 PSMain(VSOut i) : SV_Target { float4 c = Atlas.Sample(AtlasSampler, i.Uv) * i.Color * Tint; #if ALPHA_TEST clip(c.a - 0.5); #endif #if LIGHTING_LAMBERT float3 n = NormalMap.Sample(AtlasSampler, i.Uv).xyz * 2.0 - 1.0; c.rgb *= saturate(dot(n, LightDir)); #elif LIGHTING_PBR float3 n = NormalMap.Sample(AtlasSampler, i.Uv).xyz * 2.0 - 1.0; float3 l = LightDir; float d = saturate(dot(n, l)); c.rgb = c.rgb * d + pow(d, 16.0) * 0.04; #endif return c; } ``` `Program.cs`, with `RecursiveAI.OneShader` 1.2.0: ```csharp using OneShader; // A material system asks for the same few variants again and again. var cache = new ShaderCompileCache(maxEntries: 64); // thread-safe; least recently used results are dropped string[] lighting = ["NONE", "LAMBERT", "PBR"]; CompileResult Get(string light, bool alphaTest) { List> defines = [new("LIGHTING_" + light, "1")]; if (alphaTest) defines.Add(new("ALPHA_TEST", null)); // null defines it as 1 return cache.CompileFile("Shaders/Material.hlsl", new CompileOptions { Defines = defines, IncludeDirectories = ["Shaders/Common"], }); } // Three frames ask for all six variants each, from four threads at once. for (int frame = 0; frame < 3; frame++) { Parallel.ForEach(lighting.SelectMany(l => new[] { (l, false), (l, true) }), new ParallelOptions { MaxDegreeOfParallelism = 4 }, v => { if (!Get(v.Item1, v.Item2).Success) throw new Exception("compile failed"); }); Console.WriteLine($"frame {frame}: {cache.Misses} compiles, {cache.Hits} hits, {cache.Count} cached"); } // A cached result has the bytes a direct compile gives (and that call was one more hit). byte[] cached = Get("PBR", true).Find("PSMain", CompileTarget.SpirV)!.ToArray(); byte[] direct = ShaderCompiler.CompileFile("Shaders/Material.hlsl", new CompileOptions { Defines = [new("LIGHTING_PBR", "1"), new("ALPHA_TEST", null)], IncludeDirectories = ["Shaders/Common"], }).Find("PSMain", CompileTarget.SpirV)!.ToArray(); Console.WriteLine($"same bytes as a direct compile: {cached.AsSpan().SequenceEqual(direct)}"); // Editing an include is noticed: the next call for an affected variant compiles again. File.AppendAllText("Shaders/Common/Lighting.hlsli", "// edited\n"); Get("PBR", true); Console.WriteLine($"after editing Lighting.hlsli: {cache.Misses} compiles, {cache.Hits} hits"); cache.Clear(); Console.WriteLine($"after Clear(): {cache.Count} cached"); ``` ```text frame 0: 6 compiles, 0 hits, 6 cached frame 1: 6 compiles, 6 hits, 6 cached frame 2: 6 compiles, 12 hits, 6 cached same bytes as a direct compile: True after editing Lighting.hlsli: 7 compiles, 13 hits after Clear(): 0 cached ``` ## Good to know - **Concurrency.** The cache is thread-safe. Callers asking for the same key at the same time share one compile, and different keys compile in parallel. - **Bounded.** `maxEntries` (default 256) caps the number of results kept. The least recently used are dropped. - **Errors are cached too.** A failed compile is a result like any other, until its source, an include or an option changes. A cancelled call caches nothing. - **Shared results.** A cached `CompileResult` is one object handed to every caller, so treat it as read-only. Its bytes are the ones a direct `ShaderCompiler.Compile` gives. - **`Compile` and `CompileFile`.** `cache.Compile(source, path, options)` takes the text. `cache.CompileFile(path, options)` reads the file on every call, so an edit to the main file is noticed as well. For variants known at build time, [Axes in dotnet build](https://recursiveai.org/oneshader/cookbook/shader-variants-axes/) compiles them ahead of time. For the reload loop around a cache, see [Reload shaders while the app runs](https://recursiveai.org/oneshader/cookbook/hot-reload-shaders/). The cache is described in [the compiler package](https://recursiveai.org/oneshader/docs/compiler/). # Write an HLSL compute shader for Vulkan and WebGPU Write one HLSL compute shader, compile it to SPIR-V and WGSL with `dotnet build`, and get the thread-group size for the dispatch and the descriptor and bind group layouts for both APIs from C#. `Shaders/Brighten.hlsl` scales an image into a storage texture: ```hlsl cbuffer Params : register(b0) { float Gain; uint2 Size; }; Texture2D Source : register(t1); [[vk::image_format("rgba8")]] RWTexture2D Target : register(u2); [shader("compute")] [numthreads(8, 8, 1)] void CSMain(uint3 id : SV_DispatchThreadID) { if (any(id.xy >= Size)) return; Target[id.xy] = saturate(Source[id.xy] * Gain); } ``` The project references `RecursiveAI.OneShader.Build` 1.2.0 with `OneShaderTargets` set to `SpirV;Wgsl`. `Program.cs`: ```csharp using OneShader.Layouts.Vulkan; using OneShader.Layouts.WebGpu; using Compute.Shaders; // The thread-group size from [numthreads], for the dispatch. uint width = 1920, height = 1080; uint groupsX = (width + Brighten.CSMain.ThreadGroupSizeX - 1) / Brighten.CSMain.ThreadGroupSizeX; uint groupsY = (height + Brighten.CSMain.ThreadGroupSizeY - 1) / Brighten.CSMain.ThreadGroupSizeY; Console.WriteLine($"numthreads({Brighten.CSMain.ThreadGroupSizeX}, {Brighten.CSMain.ThreadGroupSizeY}, {Brighten.CSMain.ThreadGroupSizeZ}): dispatch {groupsX} x {groupsY} x 1"); // Vulkan: vkCmdDispatch(cmd, groupsX, groupsY, 1) with this layout. foreach (DescriptorSetLayoutInfo set in Brighten.Pipeline.ToVulkan().SetLayouts) foreach (DescriptorSetLayoutBindingInfo b in set.Bindings) Console.WriteLine($"Vulkan set {set.Set} binding {b.Binding}: {b.Name}, {b.DescriptorType}"); // WebGPU: pass.dispatchWorkgroups(groupsX, groupsY) with this bind group layout. if (Brighten.Pipeline.TryToWebGpu(out WebGpuPipelineLayout? webgpu, out var error)) foreach (BindGroupLayoutEntryInfo e in webgpu!.Groups[0].Entries) Console.WriteLine($"WebGPU binding {e.Binding}: {e.Name}, {e.Kind}{(e.Kind == BindingKind.StorageTexture ? $" {e.StorageFormat} {e.Access.ToWebGpuString()}" : "")}"); else Console.WriteLine(error); Console.WriteLine(System.Text.Encoding.UTF8.GetString(Brighten.CSMain.Wgsl).Split('\n').First(l => l.StartsWith("@compute"))); ``` ```text numthreads(8, 8, 1): dispatch 240 x 135 x 1 Vulkan set 0 binding 0: Params, UniformBuffer Vulkan set 0 binding 1: Source, SampledImage Vulkan set 0 binding 2: Target, StorageImage WebGPU binding 0: Params, Buffer WebGPU binding 1: Source, Texture WebGPU binding 2: Target, StorageTexture rgba8unorm write-only @compute @workgroup_size(8, 8, 1) ``` The same shader compiled with `oneshader --target spirv --target wgsl` passes `spirv-val` and naga. ## Good to know - **`ThreadGroupSizeX`, `Y` and `Z`** are constants on the generated entry-point class, taken from `[numthreads]`. Divide the work size by them, rounding up, and guard the edge in the shader as above. - **`Brighten.Pipeline`** is the file’s one compute entry point. Its `ToVulkan()` gives the descriptor set layouts and push-constant ranges, and `TryToWebGpu()` the bind group layouts. - **Storage textures on WebGPU need a format.** Without `[[vk::image_format("rgba8")]]` the build fails: `error OSH5202: The storage texture 'Target' has no texel format; WGSL needs one.` WebGPU’s access mode can’t be read from the shader, so it defaults to `write-only`. Pass `WebGpuLayoutOptions.StorageAccess` for others. A storage texture that’s both read and written has to be `r32float`, `r32uint` or `r32sint` on WGSL. - **Structured and byte-address buffers**, `groupshared` memory, barriers and 32-bit atomics work on all three targets, except compare-exchange and signed min/max on WGSL at 1.2.0 (see [InterlockedAdd and float atomics](https://recursiveai.org/oneshader/cookbook/hlsl-interlockedadd/)). Wave intrinsics need the device’s subgroup support: see [Wave intrinsics on Vulkan, Metal and WebGPU](https://recursiveai.org/oneshader/cookbook/hlsl-wave-intrinsics/). - **WebGPU doesn’t have** `Buffer`/`RWBuffer`, append and consume buffers, or atomics on textures. Each is an OSH5202 error when WGSL is selected. The full list is on [WGSL (WebGPU)](https://recursiveai.org/oneshader/docs/wgsl-backend/). To count or sum across threads, see [InterlockedAdd and float atomics](https://recursiveai.org/oneshader/cookbook/hlsl-interlockedadd/). # InterlockedAdd in HLSL, and what to do for floats Count with `InterlockedAdd` in an HLSL compute shader on Vulkan, Metal and WebGPU, and add floats atomically with a compare-exchange loop, since there is no float `InterlockedAdd`. Includes what OneShader 1.2.0 doesn’t compile. `InterlockedAdd`, `Min`, `Max`, `And`, `Or`, `Xor`, `Exchange` and `CompareExchange` work on 32-bit `int` and `uint`, in RW structured and byte-address buffers and in `groupshared` memory. Since 1.2.0 they also work on `int64_t` and `uint64_t` on SPIR-V (see below). A histogram shows both the plain counter and the float sum, `Histogram.hlsl`: ```hlsl Texture2D Image : register(t0); RWStructuredBuffer Bins : register(u1); // 256 bins RWStructuredBuffer Total : register(u2); // one float, stored as its bits // There is no float InterlockedAdd: add with a compare-exchange loop on the bits. void AtomicAddFloat(uint index, float value) { uint expected = Total[index]; for (;;) { uint desired = asuint(asfloat(expected) + value); uint original; InterlockedCompareExchange(Total[index], expected, desired, original); if (original == expected) break; expected = original; } } [shader("compute")] [numthreads(16, 16, 1)] void CSMain(uint3 id : SV_DispatchThreadID) { float3 rgb = Image[id.xy].rgb; float luma = dot(rgb, float3(0.2126, 0.7152, 0.0722)); uint bin = min(uint(luma * 256.0), 255u); InterlockedAdd(Bins[bin], 1u); // 32-bit int or uint AtomicAddFloat(0, luma); } ``` ```sh oneshader --target spirv --target msl --target wgsl Histogram.hlsl -Fo out/histogram spirv-val --target-env vulkan1.2 out/histogram.spv naga out/histogram.wgsl ``` The compile succeeds for all three targets and writes `histogram.spv`, `histogram.metal` and `histogram.wgsl`. The SPIR-V passes `spirv-val`. The Metal output uses `atomic_fetch_add_explicit` and `atomic_compare_exchange_weak_explicit`, and the WGSL output uses `atomicAdd` on `array>`. naga, however, rejects the WGSL: ```text error: Function [0] 'osh_atomic_cas_storage_u32' is invalid ┌─ histogram.wgsl:7:1 │ 7 │ ╭ fn osh_atomic_cas_storage_u32(p: ptr, read_write>, compare: u32, value: u32) -> u32 { ... = Argument 'p' at index 0 is a pointer of space Storage { access: StorageAccess(LOAD | STORE) }, which can't be passed into functions. ``` ## Limits in 1.2.0 - **Compare-exchange on WGSL.** For `InterlockedCompareExchange` and `InterlockedCompareStore` (on buffers and on `groupshared` memory), and for `InterlockedMin`/`Max` on signed `int` buffers, the WGSL that 1.2.0 writes passes a pointer to a helper function, and naga rejects that. So the float add above works on Vulkan and Metal, but not on WebGPU yet. `InterlockedAdd`, `And`, `Or`, `Xor` and `Exchange`, and `Min`/`Max` on `uint`, produce WGSL that validates. - **No float atomics.** `InterlockedAdd` on a `float` is an error that says what to do: ```text floatadd.hlsl(3,62): error OSH0314: 'InterlockedAdd' works on int, uint, int64_t and uint64_t, but the first argument is 'float'. Float atomics are not supported; reinterpret the data as uint with asuint()/asfloat(). ``` - **64-bit atomics are mostly SPIR-V.** On SPIR-V, 64-bit `Interlocked*` work on structured buffers and `groupshared` memory, and the device needs `Int64Atomics` (in the entry’s `RequiredFeatures`). Metal 2.4 has only unsigned `InterlockedMax`/`InterlockedMin` on a buffer whose result isn’t used, and WGSL has no 64-bit integers: other forms are errors (OSH5102, OSH5202) when those targets are selected. See [16-bit and 64-bit types](https://recursiveai.org/oneshader/cookbook/hlsl-16-bit-64-bit-types/). - **No atomics on textures** for Metal and WGSL (Metal needs 3.1). Use a buffer. ## Good to know - The compare-exchange loop retries until no other thread changed the value between the read and the write. Under heavy contention it’s slow, so reduce first: sum within the thread group in `groupshared` memory or with `WaveActiveSum`, and do one atomic per group or per wave. See [Wave intrinsics](https://recursiveai.org/oneshader/cookbook/hlsl-wave-intrinsics/). - Floating-point addition isn’t associative, so a sum built from atomics differs in the last bits from run to run, in any compiler. - The two-argument form `InterlockedAdd(dest, value)` discards the old value. The three-argument form returns it. For the dispatch and layouts around a compute shader, see [Write a compute shader for Vulkan and WebGPU](https://recursiveai.org/oneshader/cookbook/hlsl-compute-shader/). # HLSL wave intrinsics on Vulkan, Metal and WebGPU Use HLSL wave (subgroup) intrinsics such as `WavePrefixCountBits`, `WaveActiveSum` and `WaveReadLaneFirst` in one shader compiled to Vulkan, Metal and WebGPU, and find out which subgroup features the device has to support. OneShader 1.2.0 compiles the wave and quad intrinsics to all three targets: SPIR-V `GroupNonUniform*` instructions as DXC chooses them, Metal `simd_*` and `quad_*` functions, and WGSL’s `subgroups` feature. The supported intrinsics, on 32-bit `int`, `uint` and `float` scalars and vectors: - `WaveGetLaneCount`, `WaveGetLaneIndex`, `WaveIsFirstLane` - `WaveActiveAnyTrue`, `WaveActiveAllTrue`, `WaveActiveAllEqual`, `WaveActiveBallot`, `WaveActiveCountBits` - `WaveActiveSum`, `WaveActiveProduct`, `WaveActiveMin`, `WaveActiveMax`, `WaveActiveBitAnd`, `WaveActiveBitOr`, `WaveActiveBitXor` - `WavePrefixSum`, `WavePrefixProduct`, `WavePrefixCountBits` - `WaveReadLaneAt`, `WaveReadLaneFirst` - `QuadReadAcrossX`, `QuadReadAcrossY`, `QuadReadAcrossDiagonal`, `QuadReadLaneAt` `Shaders/Compact.hlsl` keeps the values above 0.5, with one atomic per wave instead of one per thread: ```hlsl StructuredBuffer Values : register(t0); RWStructuredBuffer Kept : register(u1); RWByteAddressBuffer Count : register(u2); [shader("compute")] [numthreads(64, 1, 1)] void CSMain(uint3 id : SV_DispatchThreadID) { float v = Values[id.x]; bool keep = v > 0.5; // One atomic per wave instead of one per thread. uint laneOffset = WavePrefixCountBits(keep); uint waveCount = WaveActiveCountBits(keep); uint waveBase = 0; if (WaveIsFirstLane()) Count.InterlockedAdd(0, waveCount, waveBase); waveBase = WaveReadLaneFirst(waveBase); if (keep) Kept[waveBase + laneOffset] = v; } ``` `Program.cs`, with `RecursiveAI.OneShader` 1.2.0: ```csharp using OneShader; using OneShader.Layouts; using OneShader.Layouts.WebGpu; CompileResult result = ShaderCompiler.CompileFile("Shaders/Compact.hlsl", new CompileOptions { Targets = [CompileTarget.SpirV, CompileTarget.Msl, CompileTarget.Wgsl], }); foreach (var d in result.Diagnostics) Console.WriteLine(d); if (!result.Success) return; foreach (ShaderBlob blob in result.Blobs) Console.WriteLine($"{blob.Target}: {blob.Data.Length} bytes"); File.WriteAllBytes("compact.spv", result.Find("CSMain", CompileTarget.SpirV)!.ToArray()); File.WriteAllText("compact.metal", result.Find("CSMain", CompileTarget.Msl)!.AsText()); File.WriteAllText("compact.wgsl", result.Find("CSMain", CompileTarget.Wgsl)!.AsText()); // What the device must support: Vulkan subgroup operations, and the WebGPU feature to request. var entry = result.Reflection!.EntryPoints.Single(); Console.WriteLine($"SubgroupOperations: {entry.SubgroupOperations} (0x{(uint)entry.SubgroupOperations:x})"); EntryLayout.TryCreate(result.Reflection!, "CSMain", out EntryLayout? cs, out _); PipelineLayout pipeline = PipelineLayout.Create(cs!); if (pipeline.TryToWebGpu(out WebGpuPipelineLayout? webgpu, out var error, new WebGpuLayoutOptions { ReadOnlyStorageBuffers = ["Values"] })) Console.WriteLine($"WebGPU requiredFeatures: [{string.Join(", ", webgpu!.RequiredFeatures ?? [])}]"); else Console.WriteLine(error); ``` ```text SpirV: 1576 bytes Msl: 1067 bytes Wgsl: 939 bytes SubgroupOperations: Basic, Ballot (0x9) WebGPU requiredFeatures: [subgroups] ``` What each target got, from the files the program wrote: | HLSL | SPIR-V | MSL | WGSL | |---|---|---|---| | `WavePrefixCountBits` | `OpGroupNonUniformBallotBitCount` `ExclusiveScan` of a ballot | `simd_prefix_exclusive_sum` | `subgroupExclusiveAdd` | | `WaveActiveCountBits` | `OpGroupNonUniformBallotBitCount` `Reduce` of a ballot | `simd_sum` | `subgroupAdd` | | `WaveIsFirstLane` | `OpGroupNonUniformElect` | `simd_is_first` | `subgroupElect` | | `WaveReadLaneFirst` | `OpGroupNonUniformBroadcastFirst` | `simd_broadcast_first` | `subgroupBroadcastFirst` | The SPIR-V passes `spirv-val` for Vulkan 1.2, and declares only the capabilities it uses (`GroupNonUniform` and `GroupNonUniformBallot` here). ## Checking the device - **Vulkan:** `SubgroupOperations` uses the values of `VkSubgroupFeatureFlagBits` (0x1 Basic, 0x8 Ballot), so compare it with `subgroupSupportedOperations` before creating the pipeline. The generated entry class and `Layout` carry the same value in build-package projects. - **WebGPU:** create the device with the features in `RequiredFeatures`. The WGSL starts with `enable subgroups;` and `diagnostic(off, subgroup_uniformity);`, because HLSL doesn’t require wave operations to be in uniform control flow. - **Metal:** SIMD-group functions need Metal 2.1 or later, which the MSL 2.4 output already requires. ## Limits in 1.2.0 - **`WaveMatch` and `WaveMultiPrefix*`** are errors, with a suggestion: ```text match.hlsl(3,61): error OSH0350: The intrinsic 'WaveMatch' is not supported: it needs partitioned subgroup operations, which no target offers portably (SPIR-V has only a vendor extension; MSL and WGSL have none). Build the groups with WaveActiveBallot and WaveReadLaneFirst in a loop instead. ``` - **Vertex shaders on WebGPU.** WebGPU’s subgroup functions exist only in compute and fragment shaders, so a vertex shader that uses one compiles for SPIR-V but not for WGSL: ```text vsw.hlsl(2,63): error OSH5202: The intrinsic 'WaveActiveSum' is not available in a vertex shader in WGSL: WebGPU's subgroup functions exist only in compute and fragment shaders. ``` On Metal, the lane index and count aren’t available in vertex shaders either. - **Checking WGSL with naga.** naga 30 doesn’t parse `enable subgroups;` yet, so it can’t validate this output. The repository checks wave shaders in a browser on WebGPU instead ([WebGPU execution parity](https://recursiveai.org/oneshader/docs/webgpu-parity/)). The instruction-level mapping per target is in [SPIR-V](https://recursiveai.org/oneshader/docs/spirv-backend/), [MSL](https://recursiveai.org/oneshader/docs/msl-backend/) and [WGSL](https://recursiveai.org/oneshader/docs/wgsl-backend/). Every intrinsic mapped to SPIR-V, Metal and WGSL: [HLSL wave intrinsics mapping table](https://recursiveai.org/oneshader/articles/hlsl-wave-intrinsics-spirv-metal-wgsl/). # The HLSL precise keyword on Vulkan, Metal and WebGPU What HLSL’s `precise` keyword does when one shader is compiled to SPIR-V, Metal Shading Language and WGSL with OneShader: `NoContraction` on Vulkan, fast-math pragmas off on Metal, and a warning on WebGPU, which can’t forbid fused multiply-adds. `precise` on a variable tells the compiler that the arithmetic feeding it must be computed exactly as written: no fusing a multiply and an add into one FMA, no reordering sums. The classic case is geometry that has to match bit for bit across draws, such as terrain patches sharing an edge. `Terrain.hlsl`: ```hlsl cbuffer Patch : register(b0) { float4x4 ViewProjection; float3 Origin; float HeightScale; float3 AxisU; float3 AxisV; }; Texture2D HeightMap : register(t1); SamplerState LinearClamp : register(s2); [shader("vertex")] float4 VSMain(float2 uv : TEXCOORD0) : SV_Position { float h = HeightMap.SampleLevel(LinearClamp, uv, 0); // Patches that share an edge must compute the same position bit for bit, // so the compiler may not fuse or reorder the arithmetic that feeds it. precise float3 world = Origin + uv.x * AxisU + uv.y * AxisV + float3(0, h * HeightScale, 0); return mul(ViewProjection, float4(world, 1)); } ``` ## Vulkan: NoContraction ```sh oneshader -T vs_6_0 -E VSMain -spirv Terrain.hlsl -Fo terrain.spv spirv-dis terrain.spv | grep -E "NoContraction|= OpF(Mul|Add)|= OpVectorTimesScalar" ``` ```text OpDecorate %41 NoContraction OpDecorate %42 NoContraction OpDecorate %48 NoContraction OpDecorate %49 NoContraction OpDecorate %54 NoContraction OpDecorate %56 NoContraction %41 = OpFMul %v3float %37 %40 %42 = OpFAdd %v3float %35 %41 %48 = OpFMul %v3float %44 %47 %49 = OpFAdd %v3float %42 %48 %54 = OpFMul %float %30 %53 %56 = OpFAdd %v3float %49 %55 ``` Every multiply and add that feeds `world` carries `NoContraction`, the SPIR-V decoration that forbids the driver from contracting it. DXC marks the same six operations for this shader (it writes the two vector-by-scalar products as `OpVectorTimesScalar`). The projection after `world` isn’t marked, because `precise` covers what feeds the variable, not what reads it. The file passes `spirv-val`. ## Metal: fast-math pragmas off ```sh oneshader -T vs_6_0 -E VSMain --target msl Terrain.hlsl -Fo terrain.metal grep -n "pragma" terrain.metal ``` ```text 50: #pragma clang fp reassociate(off) 51: #pragma clang fp contract(off) ``` Metal compiles with fast math by default. A function with precise arithmetic starts with these two pragmas, which turn off fused multiply-add and reordered sums whatever the host’s math mode. A precise division becomes `precise::divide`, since no pragma stops `a / b` becoming `a * (1 / b)`. ## WebGPU: a warning ```sh oneshader -T vs_6_0 -E VSMain --target wgsl Terrain.hlsl -Fo terrain.wgsl ``` ```text Terrain.hlsl(11,37): warning OSH5205: 'precise' arithmetic is compiled, but WGSL cannot forbid a multiply-add being fused or a sum being reordered, so the result may differ in the last bits from the other targets. Use 'fma' where the exact rounding matters, or leave WGSL out of the targets for this shader. ``` WGSL lets an implementation fuse a multiply and an add, and has no way to say otherwise. The shader compiles (naga validates the output), the warning appears once per entry point, and the result may differ from Vulkan and Metal in the last bits. With `-WX` on the command line, or `MSBuildTreatWarningsAsErrors` set to `true` in a build-package project, it stops the build instead. (C#’s `TreatWarningsAsErrors` doesn’t cover it.) ## Limits in 1.2.0 - **Matrix multiplies.** If the `precise` variable is the result of `mul` with a matrix (`precise float4 clip = mul(ViewProjection, p);`), OneShader 1.2.0 marks the scalar and vector arithmetic that feeds it but not the matrix-vector multiply itself, which DXC also marks `NoContraction`. Do the precise part in vector arithmetic, as above, or keep in mind that the driver may contract the matrix product. - **WGSL** can’t express `precise`, as above. Write `fma(a, b, c)` where you want a fused result on every target. The per-target rules are in [SPIR-V](https://recursiveai.org/oneshader/docs/spirv-backend/), [MSL](https://recursiveai.org/oneshader/docs/msl-backend/) and [WGSL](https://recursiveai.org/oneshader/docs/wgsl-backend/). What `precise` does and doesn’t guarantee on each target: [HLSL precise keyword and floating-point determinism](https://recursiveai.org/oneshader/articles/hlsl-precise-keyword-determinism/). # Convert HLSL to WGSL for WebGPU Compile an HLSL vertex, pixel or compute shader to WGSL for WebGPU with OneShader, from the command line, from C# or during `dotnet build`, check it with naga, and see what WGSL can’t express. OneShader writes WGSL from the same front end as its SPIR-V and Metal output, with the same bindings and the same buffer byte layout. The sample is `Sprite.hlsl` from [Getting started](https://recursiveai.org/oneshader/docs/getting-started/). ## From the command line ```sh oneshader --target wgsl Sprite.hlsl -Fo sprite.wgsl naga sprite.VSMain.wgsl naga sprite.PSMain.wgsl ``` Without `-T`, every `[shader("...")]` entry point is compiled, one file each: `sprite.VSMain.wgsl` and `sprite.PSMain.wgsl`. naga 30.0.1 prints `Validation successful` for both. The start of the vertex shader: ```text struct VSMain_Output { @builtin(position) VSMain_Position: vec4, @location(0) VSMain_Uv: vec2, @location(1) VSMain_Color: vec4, } @group(0) @binding(0) var Frame: array, 5>; ... @vertex fn VSMain(@location(0) in_VSMain_i_Position: vec2, @location(1) in_VSMain_i_Uv: vec2, @location(2) in_VSMain_i_Color: vec4) -> VSMain_Output { ``` - **Bindings** are `@group(set) @binding(binding)`, the same numbers as on Vulkan. - **The constant buffer is raw 16-byte rows** (`array, 5>` for the 80-byte `Frame`), read with bitcasts at the offsets HLSL packing gives. So the bytes you upload are the same for WebGPU, Vulkan and Metal, and one generated C# struct serves all three. Structured and byte-address buffers are `array` in the same way. - **The entry function** keeps the HLSL name, unless WGSL reserves it (`loop` becomes `loop_`). ## From C# or with dotnet build - At run time, add `CompileTarget.Wgsl` to `CompileOptions.Targets` and read the text with `result.Find("PSMain", CompileTarget.Wgsl)!.AsText()`. See [Compile at run time](https://recursiveai.org/oneshader/docs/getting-started/#compile-at-run-time). - In a build-package project, set `OneShaderTargets` set to `SpirV;Wgsl`. Each entry-point class then has `Wgsl` (UTF-8 bytes) and `WgslEntryPoint`, and `Pipeline.TryToWebGpu()` gives the bind group layouts with WebGPU’s enum strings. See [Write a compute shader for Vulkan and WebGPU](https://recursiveai.org/oneshader/cookbook/hlsl-compute-shader/). ## What WGSL can’t express A construct that WebGPU doesn’t have is a compile error, OSH5202, naming the target and what to do instead. For example, clip distances: ```hlsl cbuffer Frame : register(b0) { float4x4 ViewProjection; float4 ClipPlane; }; struct VSOut { float4 Position : SV_Position; float Clip : SV_ClipDistance0; }; [shader("vertex")] VSOut VSMain(float3 p : POSITION) { VSOut o; o.Position = mul(ViewProjection, float4(p, 1)); o.Clip = dot(float4(p, 1), ClipPlane); return o; } ``` ```sh oneshader --target wgsl Clip.hlsl -Fo clip.wgsl ``` ```text Clip.hlsl(2,53): error OSH5202: SV_ClipDistance and SV_CullDistance need WebGPU's optional clip-distances feature, which this backend does not use; clip in the pixel shader with clip() or discard instead. ``` The same file compiles for SPIR-V. In a build-package project the error only appears when WGSL is selected, which is why SPIR-V alone is the default. The others, from the [WGSL backend](https://recursiveai.org/oneshader/docs/wgsl-backend/) page: `CalculateLevelOfDetail`; wave intrinsics in vertex shaders; resource arrays (WebGPU has no binding arrays); `Buffer` and `RWBuffer`; append and consume buffers; atomics on textures; `double`, 64-bit integers and 16-bit integers (`half` works, as `f16`); integer and bool matrices; `PrimitiveID`, layer and viewport index; recursion; storage textures without a format, or read and written in a format other than `r32float`, `r32uint` or `r32sint`; bool stage variables; texel offsets on cube textures. ## Good to know - **Specialization constants** become `@id(N) override` declarations, set through the pipeline’s `constants`. See [Specialization constants](https://recursiveai.org/oneshader/cookbook/specialization-constants/). - **`precise`** has no WGSL equivalent: it compiles with warning OSH5205. See [The precise keyword](https://recursiveai.org/oneshader/cookbook/hlsl-precise-keyword/). - **Wave intrinsics** use WebGPU’s `subgroups` feature in compute and fragment shaders. See [Wave intrinsics](https://recursiveai.org/oneshader/cookbook/hlsl-wave-intrinsics/). - **Compare-exchange atomics** produce WGSL that naga rejects at 1.2.0. See [InterlockedAdd and float atomics](https://recursiveai.org/oneshader/cookbook/hlsl-interlockedadd/). - **Clip-space Y** points up in WebGPU and down in Vulkan, and the shader isn’t adjusted: flip it with the viewport. - **How it’s checked.** Every corpus file either validates with naga or reports one of the errors above, and the repository runs the corpus in headless Chromium on a real WebGPU adapter against DXC’s SPIR-V on Vulkan ([WebGPU execution parity](https://recursiveai.org/oneshader/docs/webgpu-parity/)). DXC with naga compared with OneShader: [HLSL to WGSL converter](https://recursiveai.org/oneshader/articles/hlsl-to-wgsl-converter/). # Compile HLSL to Metal Shading Language Compile HLSL vertex, pixel and compute shaders to Metal Shading Language with OneShader, and get from C# the function names, the buffer, texture and sampler index of every resource, and the `MTLVertexDescriptor` data. OneShader writes one MSL source file per entry point, for Metal Shading Language 2.4 or newer (macOS 12 and iOS 15), from the same front end as its SPIR-V and with the same buffer byte layout. It writes MSL directly from its own IR, without going through SPIR-V and SPIRV-Cross. ## With dotnet build The project references `RecursiveAI.OneShader.Build` 1.2.0 with `OneShaderTargets` set to `Msl` and contains the `Shaders/Sprite.hlsl` of [Getting started](https://recursiveai.org/oneshader/docs/getting-started/). `Program.cs`: ```csharp using OneShader.Layouts.Metal; using Metal.Shaders; // The MSL source and the function name to pass to newFunction(name:). File.WriteAllBytes("sprite.vs.metal", Sprite.VSMain.Msl.ToArray()); File.WriteAllBytes("sprite.ps.metal", Sprite.PSMain.Msl.ToArray()); Console.WriteLine($"vertex function {Sprite.VSMain.MslFunctionName}, fragment function {Sprite.PSMain.MslFunctionName}"); // The argument index of each resource, per stage: setVertexBuffer, setFragmentTexture, ... if (Sprite.Pipeline.TryToMetal(out MetalPipelineLayout? metal, out var error)) foreach (MetalStageLayout stage in metal!.Stages) foreach (MetalArgument arg in stage.Arguments) Console.WriteLine($"{stage.Stage}: {arg.Name} at [[{arg.Kind.ToString().ToLowerInvariant()}({arg.Index})]]"); else Console.WriteLine(error); // MTLVertexDescriptor data for the interleaved vertex buffer. MetalVertexDescriptor vd = Sprite.VSMain.Layout.Vertex!.ToMetal(Sprite.VSMain.MslVertexBufferIndex); Console.WriteLine($"vertex buffer {vd.BufferIndex}, stride {vd.Stride}, {vd.StepFunction}"); foreach (MetalVertexAttribute a in vd.Attributes) Console.WriteLine($" attributes[{a.AttributeIndex}]: {a.Format} ({(uint)a.Format}) offset {a.Offset}"); ``` ```text vertex function VSMain, fragment function PSMain Vertex: Frame at [[buffer(0)]] Pixel: Atlas at [[texture(0)]] Pixel: AtlasSampler at [[sampler(0)]] vertex buffer 30, stride 20, PerVertex attributes[0]: Float2 (29) offset 0 attributes[1]: Float2 (29) offset 8 attributes[2]: UChar4Normalized (9) offset 16 ``` The vertex function in `sprite.vs.metal`: ```text vertex VSMain_out VSMain( VSMain_in stage_in [[stage_in]], constant Frame_1_mem& Frame [[buffer(0)]]) ``` - **Argument indices** are counted per class (buffers, textures, samplers) in declaration order, so `Atlas` is texture 0 although its Vulkan binding is 1. `MetalStageLayout.TryFind(set, binding)` goes from a Vulkan-style slot to the Metal index. - **The vertex buffer** is bound at index 30 by default, away from the resource buffers. The numbers after the formats are `MTLVertexFormat` values. - **The `Frame` struct** generated in C# has the bytes `constant Frame_1_mem&` expects, the same as on Vulkan. The MSL file carries `static_assert`s on its size and offsets, compiled when `ONESHADER_CHECK_LAYOUT` is defined. ## From the command line ```sh oneshader -T ps_6_0 -E main --target msl Tint.hlsl -Fo tint.metal ``` MSL reserves `main`, so the fragment function is `main0`: ```text fragment main0_out main0( ``` `MslFunctionName` in generated code, and `MslBackend.EntryFunctionName` at run time, give the name to pass to `newFunction(name:)`. ## Limits in 1.2.0 From the [MSL backend](https://recursiveai.org/oneshader/docs/msl-backend/) page, these are compile errors on Metal: `double`; 64-bit integer vectors in buffers; 64-bit atomics other than unsigned `InterlockedMax`/`InterlockedMin` whose result isn’t used; `SV_CullDistance`; `GetDimensions` on structured buffers (Metal can’t query their length); atomics on textures; `WaveMatch` and `WaveMultiPrefix*`; the wave lane index and count in vertex shaders; integer and bool matrices; unbounded resource arrays; jumps out of several nested loops at once; comparison sampling on 1D and 3D textures; texel offsets on cube and 1D textures; overlapping `packoffset` members. With 16-bit types enabled, `half`, `short` and `ushort` are Metal’s native types, and `long` and `ulong` are native 64-bit integers. See [16-bit and 64-bit types](https://recursiveai.org/oneshader/cookbook/hlsl-16-bit-64-bit-types/). And differences to know: - **Clip-space Y** points up in Metal and down in Vulkan. The shader isn’t adjusted, so draw with a viewport of negative height, as the repository’s Metal runner does. - **Fast math** is Metal’s default, so NaN and infinity handling isn’t guaranteed unless the host turns it off. `precise` is honored either way ([The precise keyword](https://recursiveai.org/oneshader/cookbook/hlsl-precise-keyword/)). - **Function constants.** Specialization constants become `[[function_constant(N)]]`, and a function that declares any must be created with `makeFunction(name:constantValues:)`, even to get the defaults. See [Specialization constants](https://recursiveai.org/oneshader/cookbook/specialization-constants/). - **Signed integer overflow** is undefined in MSL, where HLSL wraps. - **Checking.** The samples here ran on Linux, where Apple’s compiler isn’t available, so their `.metal` files weren’t compiled with it. The repository compiles every corpus file with `xcrun metal` and runs the corpus on a Metal GPU against DXC’s SPIR-V on Vulkan. SPIRV-Cross, Metal Shader Converter and OneShader compared: [HLSL to Metal converter](https://recursiveai.org/oneshader/articles/hlsl-to-metal-converter/). # Switch from the DXC command line Replace `dxc -spirv` with `oneshader` in a Makefile, CMake or script build: which DXC flags it takes unchanged, what it adds (MSL and WGSL targets, reflection JSON, a dependency file), and which DXC options it rejects. `oneshader`, from `RecursiveAI.OneShader.Tool`, takes DXC’s core flags, so a DXC command line that emits SPIR-V for a vertex, pixel or compute shader mostly works unchanged. It needs the .NET 10 runtime: ```sh dotnet tool install -g RecursiveAI.OneShader.Tool --version 1.2.0 oneshader --version ``` ## A Makefile, before and after `shaders/post.hlsl` includes `shaders/common/color.hlsli` and uses a macro from the command line: ```hlsl #include "color.hlsli" Texture2D Scene : register(t0); SamplerState Linear : register(s1); float4 main(float2 uv : TEXCOORD0) : SV_Target { return float4(Tonemap(Scene.Sample(Linear, uv).rgb * EXPOSURE), 1); } ``` The only change to the Makefile is the compiler’s name, plus `--depfile`, so that editing the include rebuilds the shader: ```make # Was: DXC = dxc DXC = oneshader FLAGS = -spirv -O3 -HV 2021 -WX -I shaders/common -D EXPOSURE=1.5 build/post.spv: shaders/post.hlsl @mkdir -p build $(DXC) $(FLAGS) -T ps_6_0 -E main $< -Fo $@ --depfile build/post.d -include build/post.d ``` `make`, then `cat build/post.d`, then `make` again: ```text oneshader -spirv -O3 -HV 2021 -WX -I shaders/common -D EXPOSURE=1.5 -T ps_6_0 -E main shaders/post.hlsl -Fo build/post.spv --depfile build/post.d build/post.spv: /tmp/oneshader-docs/cookbook/Dxc/shaders/post.hlsl /tmp/oneshader-docs/cookbook/Dxc/shaders/common/color.hlsli make: 'build/post.spv' is up to date. ``` After `touch shaders/common/color.hlsli`, `make` compiles again. The output passes `spirv-val`, and DXC compiles the same command line (without `--depfile`) too. ## The flags From the tool’s documentation, checked against `oneshader --help` 1.2.0: | DXC flag | In oneshader | |---|---| | `-T ` | `vs_6_x`, `ps_6_x`, `cs_6_x` (also `_5_x`). The stage comes from the profile; the model number is accepted and otherwise unused. | | `-E ` | Entry point. With `-T` and no `-E` it’s `main`, as in DXC. | | `-D NAME[=VALUE]`, `-I ` | As in DXC. | | `-Fo ` | Output file. With several entry points or targets, the names are derived from it (`name..spv`, `.metal`, `.wgsl`). Without `-Fo`, a single output goes to standard output. | | `-spirv` | SPIR-V, which is also the default. (DXC’s default is DXIL, which oneshader doesn’t write.) | | `-O0`, `-Od`, `-O1`..`-O3` | Optimizer off or on (on by default). | | `-Zpr`, `-Zpc` | Row-major or column-major default matrix packing. | | `-HV 2021` | Accepted (`202x` too). Other versions are rejected. | | `-WX` | Warnings are errors. | | `-enable-16bit-types` | `half`, `min16*` and the `int16_t` family are native 16-bit types, as in DXC (since 1.2.0). See [16-bit and 64-bit types](https://recursiveai.org/oneshader/cookbook/hlsl-16-bit-64-bit-types/). | | `-fspv-target-env=vulkan1.2` | Accepted (`vulkan1.3` too). The output is valid for Vulkan 1.2 and newer. | | `-nologo` | Accepted and ignored. | What oneshader adds: | Option | Meaning | |---|---| | `--target spirv\|msl\|wgsl` | Target; repeat for several. | | `--all-entries` | Compile every `[shader("vertex"\|"pixel"\|"compute")]` entry point. Also what happens without `-T`. | | `--reflect ` | Reflection as JSON, deterministic and without machine paths. See [Read shader reflection](https://recursiveai.org/oneshader/cookbook/spirv-reflection-csharp/). | | `--depfile ` | A Make-style dependency file, written on success. | | `--no-names` | Leave source names out of the output. | Diagnostics go to standard error as `file(line,col): error OSHnnnn: message`, the form IDEs and CMake parse. The exit code is 0 for success, 1 when the shader doesn’t compile, 2 for a wrong command line. ## What it rejects Every other `-fvk-*` and `-fspv-*` option, and DXC options such as `-Zi` or `-Fh`, stop with a message that names the option (exit code 2): ```text oneshader: error: option '-Zi' is not supported: debug info is not implemented. Remove it, or use --no-names to drop names from the output. Run 'oneshader --help' for the options. oneshader: error: option '-fvk-use-dx-layout' is not supported. oneshader always uses DXC's default Vulkan layout and binding rules; use [[vk::binding]], [[vk::location]] and [[vk::offset]] in the shader instead. Run 'oneshader --help' for the options. oneshader: error: -T gs_6_0 is not a supported profile. Use vs_6_x, ps_6_x or cs_6_x; geometry, tessellation, mesh and library profiles are not supported. Run 'oneshader --help' for the options. ``` - **Layout and binding remapping** (`-fvk-b-shift`, `-fvk-use-dx-layout` and the like) isn’t supported. oneshader always applies DXC’s default Vulkan rules, so move the numbers into the shader with `[[vk::binding]]`, `[[vk::location]]` and `[[vk::offset]]`. - **DXIL**, and geometry, tessellation, mesh, ray-tracing and library profiles, aren’t supported. Keep DXC for those. ## Speed Started once per shader, as in a Makefile, the tool’s speed depends on how it’s built. The repository measures a NativeAOT build at about 2.4 times faster than DXC’s release binary per process, warm. Run as an ordinary .NET process, like the .NET tool installed above, it’s slower than DXC: about 6 times with ReadyToRun and 14 times with the JIT, because starting the runtime dominates a compile that takes 0.34 ms at the median. To get the NativeAOT binary, publish it from the repository with `dotnet publish src/OneShader.Tool -c Release -r linux-x64 -o out`. The numbers and the method are in [Performance against DXC](https://recursiveai.org/oneshader/docs/performance/). How DXC, glslang and OneShader differ on bindings and target environments: [HLSL to SPIR-V compiler compared](https://recursiveai.org/oneshader/articles/hlsl-to-spirv-compiler/).