Write an HLSL compute shader for Vulkan and WebGPU · oneshader cookbook

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<float4> Source : register(t1);
[[vk::image_format("rgba8")]] RWTexture2D<float4> 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:

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). Wave intrinsics need the device’s subgroup support: see Wave intrinsics on Vulkan, Metal and WebGPU.
WebGPU doesn’t have Buffer<T>/RWBuffer<T>, append and consume buffers, or atomics on textures. Each is an OSH5202 error when WGSL is selected. The full list is on WGSL (WebGPU).

To count or sum across threads, see InterlockedAdd and float atomics.