---
title: "Write an HLSL compute shader for Vulkan and WebGPU"
url: https://recursiveai.org/oneshader/cookbook/hlsl-compute-shader/
description: "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…"
version: "1.2.0"
---

# 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<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`:

```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<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)](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/).