From 6ac590e3878720bd6599984ac6eadf8526781613 Mon Sep 17 00:00:00 2001 From: Martin Valigursky Date: Tue, 2 Jun 2026 09:45:40 +0100 Subject: [PATCH 1/3] feat(webgpu): automatic bind-group reflection for compute shaders WebGPU compute shaders can now use the same simplified WGSL syntax as vertex/fragment shaders. The engine reflects resources from the shader source and builds the bind group automatically, so a hand-written computeBindGroupFormat is no longer required. - Reflect loose uniforms (into a generated uniform buffer), textures + samplers, storage buffers, and storage textures declared with simplified syntax. - computeBindGroupFormat / computeUniformBufferFormats are now optional. Reflected resources go in their own bind group (group 1 when a caller format is supplied, otherwise group 0); explicitly-bound resources are left untouched. - Generalize compute pipeline / bind-group execution to support multiple bind groups. - Fix storage-buffer reflection regex to accept read_write. - Convert edge-detect and particles examples to the simplified syntax. - Add unit tests for the compute reflection paths. Fixes #7689 --- .../compute/edge-detect.compute-shader.wgsl | 8 +- .../examples/compute/edge-detect.example.mjs | 13 +- .../examples/compute/particles.example.mjs | 26 +- .../compute/particles.shader-simulation.wgsl | 25 +- .../webgpu/webgpu-compute-pipeline.js | 26 +- .../graphics/webgpu/webgpu-compute.js | 86 ++++-- .../webgpu/webgpu-shader-processor-wgsl.js | 158 +++++++++- src/platform/graphics/webgpu/webgpu-shader.js | 65 ++++- ...gpu-shader-processor-wgsl-compute.test.mjs | 276 ++++++++++++++++++ 9 files changed, 587 insertions(+), 96 deletions(-) create mode 100644 test/platform/graphics/webgpu/webgpu-shader-processor-wgsl-compute.test.mjs diff --git a/examples/src/examples/compute/edge-detect.compute-shader.wgsl b/examples/src/examples/compute/edge-detect.compute-shader.wgsl index 144b0644e47..39c01da290b 100644 --- a/examples/src/examples/compute/edge-detect.compute-shader.wgsl +++ b/examples/src/examples/compute/edge-detect.compute-shader.wgsl @@ -1,9 +1,11 @@ // Include half-precision type aliases (resolves to f16 when supported, f32 otherwise) #include "halfTypesCS" -@group(0) @binding(0) var inputTexture: texture_2d; -@group(0) @binding(1) var inputTexture_sampler: sampler; -@group(0) @binding(2) var outputTexture: texture_storage_2d; +// Simplified-syntax declarations (no @group/@binding) - the engine reflects these into a bind +// group automatically, so the example does not provide a computeBindGroupFormat. +var inputTexture: texture_2d; +var inputTexture_sampler: sampler; +var outputTexture: texture_storage_2d; @compute @workgroup_size(8, 8, 1) fn main(@builtin(global_invocation_id) global_id : vec3u) { diff --git a/examples/src/examples/compute/edge-detect.example.mjs b/examples/src/examples/compute/edge-detect.example.mjs index a52bd6f2ce2..df774920b47 100644 --- a/examples/src/examples/compute/edge-detect.example.mjs +++ b/examples/src/examples/compute/edge-detect.example.mjs @@ -173,18 +173,13 @@ assetListLoader.load(() => { const createComputeShader = () => { if (!device.supportsCompute) return null; + // No computeBindGroupFormat is provided - the input texture (+ sampler) and the output + // storage texture use the simplified WGSL syntax and are reflected automatically by the + // engine from the shader source. return new pc.Shader(device, { name: 'EdgeDetect-Shader', shaderLanguage: pc.SHADERLANGUAGE_WGSL, - cshader: computeShaderWgsl, - - // Format of a bind group for the compute shader - computeBindGroupFormat: new pc.BindGroupFormat(device, [ - // Input texture with sampler (sampler takes binding slot+1 automatically) - new pc.BindTextureFormat('inputTexture', pc.SHADERSTAGE_COMPUTE, undefined, undefined, true), - // Output storage texture - new pc.BindStorageTextureFormat('outputTexture', pc.PIXELFORMAT_RGBA8, pc.TEXTUREDIMENSION_2D) - ]) + cshader: computeShaderWgsl }); }; diff --git a/examples/src/examples/compute/particles.example.mjs b/examples/src/examples/compute/particles.example.mjs index 2b1db01f789..a205c9a35f6 100644 --- a/examples/src/examples/compute/particles.example.mjs +++ b/examples/src/examples/compute/particles.example.mjs @@ -84,31 +84,15 @@ assetListLoader.load(() => { const numParticles = 1024 * 1024; - // a compute shader that will simulate the particles stored in a storage buffer + // a compute shader that will simulate the particles stored in a storage buffer. No bind group + // or uniform buffer formats are provided - the loose uniforms (count, dt, sphereCount) and the + // storage buffers (particles, spheres) use the simplified WGSL syntax and are reflected + // automatically by the engine from the shader source. const shader = device.supportsCompute ? new pc.Shader(device, { name: 'SimulationShader', shaderLanguage: pc.SHADERLANGUAGE_WGSL, - cshader: shaderSharedWgsl + shaderSimulationWgsl, - - // format of a uniform buffer used by the compute shader - computeUniformBufferFormats: { - ub: new pc.UniformBufferFormat(device, [ - new pc.UniformFormat('count', pc.UNIFORMTYPE_UINT), - new pc.UniformFormat('dt', pc.UNIFORMTYPE_FLOAT), - new pc.UniformFormat('sphereCount', pc.UNIFORMTYPE_UINT) - ]) - }, - - // format of a bind group, providing resources for the compute shader - computeBindGroupFormat: new pc.BindGroupFormat(device, [ - // a uniform buffer we provided the format for - new pc.BindUniformBufferFormat('ub', pc.SHADERSTAGE_COMPUTE), - // particle storage buffer - new pc.BindStorageBufferFormat('particles', pc.SHADERSTAGE_COMPUTE), - // rad only collision spheres - new pc.BindStorageBufferFormat('spheres', pc.SHADERSTAGE_COMPUTE, true) - ]) + cshader: shaderSharedWgsl + shaderSimulationWgsl }) : null; diff --git a/examples/src/examples/compute/particles.shader-simulation.wgsl b/examples/src/examples/compute/particles.shader-simulation.wgsl index e33564d5e29..bf16daa5000 100644 --- a/examples/src/examples/compute/particles.shader-simulation.wgsl +++ b/examples/src/examples/compute/particles.shader-simulation.wgsl @@ -1,30 +1,29 @@ -// uniform buffer for the compute shader -struct ub_compute { - count: u32, // number of particles - dt: f32, // delta time - sphereCount: u32 // number of spheres -} - // sphere struct used for the colliders struct Sphere { center: vec3, radius: f32 } -@group(0) @binding(0) var ubCompute : ub_compute; -@group(0) @binding(1) var particles: array; -@group(0) @binding(2) var spheres: array; +// Simplified-syntax resources (no @group/@binding) - the engine reflects these into a bind group +// automatically, so the example does not provide computeBindGroupFormat / computeUniformBufferFormats. +// The loose uniforms are collapsed into a single generated uniform buffer. +uniform count: u32; // number of particles +uniform dt: f32; // delta time +uniform sphereCount: u32; // number of spheres + +var particles: array; +var spheres: array; @compute @workgroup_size(64) fn main(@builtin(global_invocation_id) global_invocation_id: vec3u) { // particle index - ignore if out of bounds (as they get batched into groups of 64) let index = global_invocation_id.x * 1024 + global_invocation_id.y; - if (index >= ubCompute.count) { return; } + if (index >= uniform.count) { return; } // update times var particle = particles[index]; - particle.collisionTime += ubCompute.dt; + particle.collisionTime += uniform.dt; // if particle gets too far, reset it to its original position / velocity var distance = length(particle.position); @@ -41,7 +40,7 @@ fn main(@builtin(global_invocation_id) global_invocation_id: vec3u) { var next = particle.position + delta; // handle collisions with spheres - for (var i = 0u; i < ubCompute.sphereCount; i++) { + for (var i = 0u; i < uniform.sphereCount; i++) { var center = spheres[i].center; var radius = spheres[i].radius; diff --git a/src/platform/graphics/webgpu/webgpu-compute-pipeline.js b/src/platform/graphics/webgpu/webgpu-compute-pipeline.js index 6259d003a22..6f5db71d378 100644 --- a/src/platform/graphics/webgpu/webgpu-compute-pipeline.js +++ b/src/platform/graphics/webgpu/webgpu-compute-pipeline.js @@ -29,7 +29,8 @@ class CacheEntry { } class WebgpuComputePipeline extends WebgpuPipeline { - lookupHashes = new Uint32Array(2); + // shader compute key + up to 2 bind group format keys (caller group 0 + reflected group) + lookupHashes = new Uint32Array(3); /** * The cache of compute pipelines @@ -38,12 +39,22 @@ class WebgpuComputePipeline extends WebgpuPipeline { */ cache = new Map(); - get(shader, bindGroupFormat) { + /** + * @param {import('../shader.js').Shader} shader - The compute shader. + * @param {import('../bind-group-format.js').BindGroupFormat[]} bindGroupFormats - The bind group + * formats, in bind group index order (dense, no gaps). + * @returns {GPUComputePipeline} - The compute pipeline. + */ + get(shader, bindGroupFormats) { + + Debug.assert(bindGroupFormats.length <= 2); - // unique hash for the pipeline + // unique hash for the pipeline - shader key followed by each bind group format key (0 for + // an absent group). All slots are written, so no need to clear stale values from reuse. const lookupHashes = this.lookupHashes; lookupHashes[0] = shader.impl.computeKey; - lookupHashes[1] = bindGroupFormat.impl.key; + lookupHashes[1] = bindGroupFormats[0] ? bindGroupFormats[0].impl.key : 0; + lookupHashes[2] = bindGroupFormats[1] ? bindGroupFormats[1].impl.key : 0; const hash = hash32Fnv1a(lookupHashes); // Check cache @@ -58,8 +69,11 @@ class WebgpuComputePipeline extends WebgpuPipeline { } } - // Cache miss - create new pipeline - const pipelineLayout = this.getPipelineLayout([bindGroupFormat.impl]); + // Cache miss - create new pipeline. Build the impl array explicitly (at most 2 groups). + const impls = []; + if (bindGroupFormats[0]) impls.push(bindGroupFormats[0].impl); + if (bindGroupFormats[1]) impls.push(bindGroupFormats[1].impl); + const pipelineLayout = this.getPipelineLayout(impls); const cacheEntry = new CacheEntry(); cacheEntry.hashes = new Uint32Array(lookupHashes); cacheEntry.pipeline = this.create(shader, pipelineLayout); diff --git a/src/platform/graphics/webgpu/webgpu-compute.js b/src/platform/graphics/webgpu/webgpu-compute.js index af077615a6a..28b68b44daa 100644 --- a/src/platform/graphics/webgpu/webgpu-compute.js +++ b/src/platform/graphics/webgpu/webgpu-compute.js @@ -15,8 +15,14 @@ class WebgpuCompute { /** @type {UniformBuffer[]} */ uniformBuffers = []; - /** @type {BindGroup} */ - bindGroup = null; + /** + * Bind groups, indexed by bind group index. A caller-provided format occupies group 0; + * auto-reflected resources occupy their own group (0 when no caller format, otherwise 1). + * The array is dense (no gaps), as required by WebGPU pipeline layouts. + * + * @type {BindGroup[]} + */ + bindGroups = []; constructor(compute) { this.compute = compute; @@ -25,27 +31,53 @@ class WebgpuCompute { DebugGraphics.pushGpuMarker(device, `Compute:${compute.name}`); - // create bind group - const { computeBindGroupFormat, computeUniformBufferFormats } = shader.impl; - Debug.assert(computeBindGroupFormat, 'Compute shader does not have computeBindGroupFormat specified', shader); - - // this.bindGroup = new BindGroup(device, computeBindGroupFormat, this.uniformBuffer); - this.bindGroup = new BindGroup(device, computeBindGroupFormat); - DebugHelper.setName(this.bindGroup, `Compute-BindGroup_${this.bindGroup.id}`); - - if (computeUniformBufferFormats) { - for (const name in computeUniformBufferFormats) { - if (computeUniformBufferFormats.hasOwnProperty(name)) { - // TODO: investigate implications of using a non-persistent uniform buffer - const ub = new UniformBuffer(device, computeUniformBufferFormats[name], true); - this.uniformBuffers.push(ub); - this.bindGroup.setUniformBuffer(name, ub); + const { + computeBindGroupFormat, computeUniformBufferFormats, + computeReflectedBindGroupFormat, computeReflectedUniformBufferFormat, + computeReflectedGroupIndex + } = shader.impl; + + // ordered, gapless array of bind group formats (array index === bind group index) + const formats = []; + + // group 0: caller-provided resources (if any) + if (computeBindGroupFormat) { + const bindGroup = new BindGroup(device, computeBindGroupFormat); + DebugHelper.setName(bindGroup, `Compute-BindGroup_${bindGroup.id}`); + + if (computeUniformBufferFormats) { + for (const name in computeUniformBufferFormats) { + if (computeUniformBufferFormats.hasOwnProperty(name)) { + // TODO: investigate implications of using a non-persistent uniform buffer + const ub = new UniformBuffer(device, computeUniformBufferFormats[name], true); + this.uniformBuffers.push(ub); + bindGroup.setUniformBuffer(name, ub); + } } } + + formats[0] = computeBindGroupFormat; + this.bindGroups[0] = bindGroup; + } + + // auto-reflected resources, at their own bind group (0 when no caller format, otherwise 1) + if (computeReflectedBindGroupFormat) { + const reflectedBindGroup = new BindGroup(device, computeReflectedBindGroupFormat); + DebugHelper.setName(reflectedBindGroup, `Compute-ReflectedBindGroup_${reflectedBindGroup.id}`); + + if (computeReflectedUniformBufferFormat) { + // matches the generated 'ub_compute' uniform buffer (see WebgpuShaderProcessorWGSL.runCompute) + const ub = new UniformBuffer(device, computeReflectedUniformBufferFormat, true); + this.uniformBuffers.push(ub); + reflectedBindGroup.setUniformBuffer('ub_compute', ub); + } + + formats[computeReflectedGroupIndex] = computeReflectedBindGroupFormat; + this.bindGroups[computeReflectedGroupIndex] = reflectedBindGroup; } // pipeline - this.pipeline = device.computePipeline.get(shader, computeBindGroupFormat); + this.pipeline = device.computePipeline.get(shader, formats); DebugGraphics.popGpuMarker(device); } @@ -55,23 +87,27 @@ class WebgpuCompute { this.uniformBuffers.forEach(ub => ub.destroy()); this.uniformBuffers.length = 0; - this.bindGroup.destroy(); - this.bindGroup = null; + this.bindGroups.forEach(bindGroup => bindGroup.destroy()); + this.bindGroups.length = 0; } updateBindGroup() { // bind group data - const { bindGroup } = this; - bindGroup.updateUniformBuffers(); - bindGroup.update(); + for (let i = 0; i < this.bindGroups.length; i++) { + const bindGroup = this.bindGroups[i]; + bindGroup.updateUniformBuffers(); + bindGroup.update(); + } } dispatch(x, y, z) { - // bind group + // bind groups const device = this.compute.device; - device.setBindGroup(0, this.bindGroup); + for (let i = 0; i < this.bindGroups.length; i++) { + device.setBindGroup(i, this.bindGroups[i]); + } // compute pipeline const passEncoder = device.passEncoder; diff --git a/src/platform/graphics/webgpu/webgpu-shader-processor-wgsl.js b/src/platform/graphics/webgpu/webgpu-shader-processor-wgsl.js index b3315aeb1ea..2f6c6a43e5e 100644 --- a/src/platform/graphics/webgpu/webgpu-shader-processor-wgsl.js +++ b/src/platform/graphics/webgpu/webgpu-shader-processor-wgsl.js @@ -1,7 +1,7 @@ import { Debug } from '../../../core/debug.js'; import { BINDGROUP_MESH, semanticToLocation, - SHADERSTAGE_VERTEX, SHADERSTAGE_FRAGMENT, + SHADERSTAGE_VERTEX, SHADERSTAGE_FRAGMENT, SHADERSTAGE_COMPUTE, SAMPLETYPE_FLOAT, TEXTUREDIMENSION_2D, TEXTUREDIMENSION_2D_ARRAY, TEXTUREDIMENSION_CUBE, TEXTUREDIMENSION_3D, TEXTUREDIMENSION_1D, TEXTUREDIMENSION_CUBE_ARRAY, @@ -15,7 +15,8 @@ import { TYPE_FLOAT32, TYPE_FLOAT16, TYPE_INT8, TYPE_INT16, TYPE_INT32 } from '../constants.js'; import { UniformFormat, UniformBufferFormat } from '../uniform-buffer-format.js'; -import { BindGroupFormat, BindStorageBufferFormat, BindTextureFormat } from '../bind-group-format.js'; +import { BindGroupFormat, BindStorageBufferFormat, BindStorageTextureFormat, BindTextureFormat, BindUniformBufferFormat } from '../bind-group-format.js'; +import { gpuTextureFormats } from './constants.js'; /** * @import { GraphicsDevice } from '../graphics-device.js' @@ -173,6 +174,18 @@ const getTextureDeclarationType = (viewDimension, sampleType) => { return `${baseTypeString}<${coreFormatString}>`; }; +// reverse of gpuTextureFormats: WGSL/GPU storage format string -> PIXELFORMAT. Built once, used to +// reflect storage texture declarations. Several PIXELFORMATs can map to the same string (e.g. RGB8 +// and RGBA8 both -> 'rgba8unorm'); first-wins is fine as the chosen PIXELFORMAT only needs to +// round-trip back to the same string via gpuTextureFormats[...], which is what the bind group +// layout (and WebGPU validation) uses. +const gpuFormatToPixelFormat = new Map(); +gpuTextureFormats.forEach((str, pixelFormat) => { + if (str && !gpuFormatToPixelFormat.has(str)) { + gpuFormatToPixelFormat.set(str, pixelFormat); + } +}); + const wrappedArrayTypes = { 'f32': 'WrappedF32', 'i32': 'WrappedI32', @@ -253,7 +266,7 @@ const TEXTURE_REGEX = /^\s*var\s+(\w+)\s*:\s*(texture_\w+)(?:<(\w+)>)?;\s*$/; // eslint-disable-next-line const STORAGE_TEXTURE_REGEX = /^\s*var\s+([\w\d_]+)\s*:\s*(texture_storage_2d|texture_storage_2d_array)<([\w\d_]+),\s*(\w+)>\s*;\s*$/; // eslint-disable-next-line -const STORAGE_BUFFER_REGEX = /^\s*var\s*\s*([\w\d_]+)\s*:\s*(.*)\s*;\s*$/; +const STORAGE_BUFFER_REGEX = /^\s*var\s*\s*([\w\d_]+)\s*:\s*(.*)\s*;\s*$/; // eslint-disable-next-line const EXTERNAL_TEXTURE_REGEX = /^\s*var\s+([\w\d_]+)\s*:\s*texture_external;\s*$/; // eslint-disable-next-line @@ -445,6 +458,79 @@ class WebgpuShaderProcessorWGSL { }; } + /** + * Process a compute shader: reflect its simplified-syntax declarations (loose `uniform`s, + * textures/samplers, storage buffers) into a single bind group at `reflectedGroupIndex`, leaving + * any explicitly-bound (`@group/@binding`) declarations untouched. The loose uniforms are + * collapsed into one generated uniform buffer (`ub_compute`) placed inside that same group. + * + * @param {GraphicsDevice} device - The graphics device. + * @param {string} source - The fully-preprocessed compute shader source (includes/defines resolved). + * @param {object} shaderDefinition - The shader definition. + * @param {Shader} shader - The shader. + * @param {number} reflectedGroupIndex - The bind group index the reflected resources are placed + * in (0 when no caller format is supplied, otherwise 1). + * @returns {object} - `{ cshader, computeBindGroupFormat, computeUniformBufferFormat }`. The + * formats are null when there is nothing to reflect (strictly additive - behavior is then + * identical to a fully hand-authored compute shader). + */ + static runCompute(device, source, shaderDefinition, shader, reflectedGroupIndex) { + + // pull simplified-syntax declarations out of the source (explicit @group/@binding lines, + // which start with '@' rather than 'var'/'uniform', are not matched and pass through) + const extracted = WebgpuShaderProcessorWGSL.extract(source); + + // parse loose uniforms - all of them go into the single generated compute uniform buffer + const parsedUniforms = extracted.uniforms.map(line => new UniformLine(line, shader)); + const meshUniforms = []; + parsedUniforms.forEach((uniform) => { + uniform.ubName = 'ub_compute'; + const uniformType = uniformTypeToNameMapWGSL.get(uniform.type); + Debug.assert(uniformType !== undefined, `Uniform type ${uniform.type} is not recognised on line [${uniform.line}]`); + meshUniforms.push(new UniformFormat(uniform.name, uniformType, uniform.arraySize)); + }); + // do not synthesize a dummy uniform when empty - reflection must stay strictly additive + const computeUniformBufferFormat = meshUniforms.length > 0 ? new UniformBufferFormat(device, meshUniforms) : null; + + // parse resource lines (no vertex/fragment merge for compute) + const parsedResources = WebgpuShaderProcessorWGSL.mergeResources(extracted.resources, [], shader); + const resourceFormats = WebgpuShaderProcessorWGSL.buildResourceFormats(parsedResources, SHADERSTAGE_COMPUTE, shader); + + // the generated uniform buffer is a binding inside the same reflected group, appended last + const ubBindFormat = computeUniformBufferFormat ? new BindUniformBufferFormat('ub_compute', SHADERSTAGE_COMPUTE) : null; + const allFormats = ubBindFormat ? [...resourceFormats, ubBindFormat] : resourceFormats; + + // when there is nothing to reflect, leave the source and bindings exactly as the caller + // provided them (strictly additive); otherwise build the single reflected bind group format + // (this assigns the slots), generate the declarations and inject them into the source + let cshader = source; + let computeBindGroupFormat = null; + if (allFormats.length > 0) { + + computeBindGroupFormat = new BindGroupFormat(device, allFormats); + + // generate declarations using the assigned slots + let code = WebgpuShaderProcessorWGSL.getTextureShaderDeclaration(computeBindGroupFormat, reflectedGroupIndex); + if (computeUniformBufferFormat) { + code += WebgpuShaderProcessorWGSL.getUniformShaderDeclaration(computeUniformBufferFormat, reflectedGroupIndex, ubBindFormat.slot, 'compute'); + } + + // rewrite `uniform.x` references to `ub_compute.x` + const src = WebgpuShaderProcessorWGSL.renameUniformAccess(extracted.src, parsedUniforms); + + // insert the generated declarations at the marker (or prepend if there was no marker) + cshader = src.includes(MARKER) ? src.replace(MARKER, code) : `${code}\n${src}`; + } + + // computeUniformBufferFormat is already null unless loose uniforms were reflected (in which + // case computeBindGroupFormat is non-null too), so the result stays consistent + return { + cshader, + computeBindGroupFormat, + computeUniformBufferFormat + }; + } + // Extract required information from the shader source code. static extract(src) { // collected data @@ -612,10 +698,21 @@ class WebgpuShaderProcessorWGSL { return resources; } - static processResources(device, resources, processingOptions, shader) { + /** + * Converts parsed resource lines (textures, samplers, storage buffers) into an array of bind + * formats. Shared by the vertex/fragment path ({@link processResources}) and the compute path + * ({@link runCompute}); only the shader-stage visibility differs. + * + * @param {Array} resources - The parsed resource lines. + * @param {number} visibility - Shader stage visibility bit-flags for the created formats. + * @param {Shader} shader - The shader (for error reporting). + * @returns {Array} - The bind formats, in declaration + * order (a texture with a sampler consumes the following sampler line). + * @private + */ + static buildResourceFormats(resources, visibility, shader) { - // build mesh bind group format - this contains the textures, but not the uniform buffer as that is a separate binding - const textureFormats = []; + const formats = []; for (let i = 0; i < resources.length; i++) { const resource = resources[i]; @@ -631,7 +728,7 @@ class WebgpuShaderProcessorWGSL { const dimension = resource.textureDimension; // TODO: we could optimize visibility to only stages that use any of the data - textureFormats.push(new BindTextureFormat(resource.name, SHADERSTAGE_VERTEX | SHADERSTAGE_FRAGMENT, dimension, sampleType, hasSampler, hasSampler ? sampler.name : null)); + formats.push(new BindTextureFormat(resource.name, visibility, dimension, sampleType, hasSampler, hasSampler ? sampler.name : null)); // following sampler was already handled if (hasSampler) i++; @@ -640,28 +737,47 @@ class WebgpuShaderProcessorWGSL { if (resource.isStorageBuffer) { const readOnly = resource.accessMode !== 'read_write'; - const bufferFormat = new BindStorageBufferFormat(resource.name, SHADERSTAGE_VERTEX | SHADERSTAGE_FRAGMENT, readOnly); + const bufferFormat = new BindStorageBufferFormat(resource.name, visibility, readOnly); bufferFormat.format = resource.type; - textureFormats.push(bufferFormat); + formats.push(bufferFormat); + } + + if (resource.isStorageTexture) { + + // storage textures are compute-only (BindStorageTextureFormat hardcodes + // SHADERSTAGE_COMPUTE); the `visibility` param does not apply here + const dimension = resource.textureType === 'texture_storage_2d_array' ? TEXTUREDIMENSION_2D_ARRAY : TEXTUREDIMENSION_2D; + const pixelFormat = gpuFormatToPixelFormat.get(resource.format); + Debug.assert(pixelFormat !== undefined, `Unsupported storage texture format '${resource.format}' on line [${resource.originalLine}]`); + const write = resource.access === 'write' || resource.access === 'read_write'; + const read = resource.access === 'read' || resource.access === 'read_write'; + formats.push(new BindStorageTextureFormat(resource.name, pixelFormat, dimension, write, read)); } Debug.assert(!resource.isSampler, `Sampler uniform needs to follow a texture uniform, but does not on line [${resource.originalLine}]`); - Debug.assert(!resource.isStorageTexture, 'TODO: add support for storage textures here'); Debug.assert(!resource.externalTexture, 'TODO: add support for external textures here'); } + return formats; + } + + static processResources(device, resources, processingOptions, shader, visibility = SHADERSTAGE_VERTEX | SHADERSTAGE_FRAGMENT, bindGroupIndex = BINDGROUP_MESH) { + + // build mesh bind group format - this contains the textures, but not the uniform buffer as that is a separate binding + const textureFormats = WebgpuShaderProcessorWGSL.buildResourceFormats(resources, visibility, shader); + const meshBindGroupFormat = new BindGroupFormat(device, textureFormats); // generate code for textures let code = ''; - processingOptions.bindGroupFormats.forEach((format, bindGroupIndex) => { + processingOptions?.bindGroupFormats?.forEach((format, index) => { if (format) { - code += WebgpuShaderProcessorWGSL.getTextureShaderDeclaration(format, bindGroupIndex); + code += WebgpuShaderProcessorWGSL.getTextureShaderDeclaration(format, index); } }); // and also for generated mesh format - code += WebgpuShaderProcessorWGSL.getTextureShaderDeclaration(meshBindGroupFormat, BINDGROUP_MESH); + code += WebgpuShaderProcessorWGSL.getTextureShaderDeclaration(meshBindGroupFormat, bindGroupIndex); return { code, @@ -679,12 +795,14 @@ class WebgpuShaderProcessorWGSL { * @param {UniformBufferFormat} ubFormat - Format of the uniform buffer. * @param {number} bindGroup - The bind group index. * @param {number} bindIndex - The bind index. + * @param {string} [name] - The name used for the struct and uniform buffer variable + * (`struct_ub_` / `ub_`). Defaults to the bind group name. The compute path passes + * an explicit name as its reflected group index does not map to a meaningful bindGroupNames entry. * @returns {string} - The shader code for the uniform buffer. * @private */ - static getUniformShaderDeclaration(ubFormat, bindGroup, bindIndex) { + static getUniformShaderDeclaration(ubFormat, bindGroup, bindIndex, name = bindGroupNames[bindGroup]) { - const name = bindGroupNames[bindGroup]; const structName = `struct_ub_${name}`; let code = `struct ${structName} {\n`; @@ -749,7 +867,15 @@ class WebgpuShaderProcessorWGSL { }); - Debug.assert(format.storageTextureFormats.length === 0, 'Implement support for storage textures here'); + format.storageTextureFormats.forEach((format) => { + + const storageType = format.textureDimension === TEXTUREDIMENSION_2D_ARRAY ? 'texture_storage_2d_array' : 'texture_storage_2d'; + const fmtString = gpuTextureFormats[format.format]; + const access = format.read ? (format.write ? 'read_write' : 'read') : 'write'; + code += `@group(${bindGroup}) @binding(${format.slot}) var ${format.name}: ${storageType}<${fmtString}, ${access}>;\n`; + + }); + // TODO: also add external texture support here return code; diff --git a/src/platform/graphics/webgpu/webgpu-shader.js b/src/platform/graphics/webgpu/webgpu-shader.js index 0a5da98553a..4330838aa3e 100644 --- a/src/platform/graphics/webgpu/webgpu-shader.js +++ b/src/platform/graphics/webgpu/webgpu-shader.js @@ -49,6 +49,36 @@ class WebgpuShader { */ _computeKey; + /** + * Caller-provided bind group format (compute, group 0), or null if none was supplied. + * + * @type {import('../bind-group-format.js').BindGroupFormat|null} + */ + computeBindGroupFormat = null; + + /** + * Bind group format for resources auto-reflected from the compute shader source. Lives at + * {@link computeReflectedGroupIndex}. Null when there is nothing to reflect. + * + * @type {import('../bind-group-format.js').BindGroupFormat|null} + */ + computeReflectedBindGroupFormat = null; + + /** + * Generated uniform buffer format holding the reflected loose uniforms, bound inside the + * reflected bind group. Null when the shader declares no loose uniforms. + * + * @type {import('../uniform-buffer-format.js').UniformBufferFormat|null} + */ + computeReflectedUniformBufferFormat = null; + + /** + * Bind group index of the reflected resources (0 when no caller format, otherwise 1). + * + * @type {number} + */ + computeReflectedGroupIndex = 0; + /** * Name of the vertex entry point function. */ @@ -78,13 +108,12 @@ class WebgpuShader { if (definition.cshader) { - this._computeCode = definition.cshader ?? null; - this.computeUniformBufferFormats = definition.computeUniformBufferFormats; - this.computeBindGroupFormat = definition.computeBindGroupFormat; if (definition.computeEntryPoint) { this.computeEntryPoint = definition.computeEntryPoint; } + this.processComputeWGSL(); + } else { this.vertexEntryPoint = 'vertexMain'; @@ -181,6 +210,36 @@ class WebgpuShader { shader.attributes = processed.attributes; } + processComputeWGSL() { + const shader = this.shader; + const definition = shader.definition; + + // a caller-provided bind group format occupies group 0; otherwise reflected resources + // start at group 0 (WebGPU pipeline layouts cannot have gaps) + const callerBindGroupFormat = definition.computeBindGroupFormat ?? null; + const reflectedGroupIndex = callerBindGroupFormat ? 1 : 0; + + // reflect simplified-syntax declarations into a separate bind group, leaving any + // explicitly-bound resources (and the caller-provided format) untouched + const processed = WebgpuShaderProcessorWGSL.runCompute(shader.device, definition.cshader, definition, shader, reflectedGroupIndex); + + // keep reference to processed shader in debug mode + Debug.call(() => { + this.processed = processed; + }); + + this._computeCode = processed.cshader; + + // caller-provided (group 0) resources + this.computeBindGroupFormat = callerBindGroupFormat; + this.computeUniformBufferFormats = definition.computeUniformBufferFormats; + + // reflected (engine-managed) resources and their generated uniform buffer + this.computeReflectedGroupIndex = reflectedGroupIndex; + this.computeReflectedBindGroupFormat = processed.computeBindGroupFormat; + this.computeReflectedUniformBufferFormat = processed.computeUniformBufferFormat; + } + processWGSL() { const shader = this.shader; diff --git a/test/platform/graphics/webgpu/webgpu-shader-processor-wgsl-compute.test.mjs b/test/platform/graphics/webgpu/webgpu-shader-processor-wgsl-compute.test.mjs new file mode 100644 index 00000000000..dd6f88cf9df --- /dev/null +++ b/test/platform/graphics/webgpu/webgpu-shader-processor-wgsl-compute.test.mjs @@ -0,0 +1,276 @@ +import { expect } from 'chai'; + +import { + SHADERSTAGE_COMPUTE, + SAMPLETYPE_FLOAT, SAMPLETYPE_UNFILTERABLE_FLOAT +} from '../../../../src/platform/graphics/constants.js'; +import { ScopeSpace } from '../../../../src/platform/graphics/scope-space.js'; +import { WebgpuShaderProcessorWGSL } from '../../../../src/platform/graphics/webgpu/webgpu-shader-processor-wgsl.js'; + +// Minimal device mock - BindGroupFormat / UniformBufferFormat only need scope.resolve and +// createBindGroupFormatImpl. A real WebGPU device is not available in the headless test runner. +function createMockDevice() { + let implKey = 0; + return { + scope: new ScopeSpace('test'), + createBindGroupFormatImpl() { + return { key: implKey++, destroy() {} }; + } + }; +} + +function run(source, reflectedGroupIndex = 1) { + const device = createMockDevice(); + const shader = { failed: false, name: 'test' }; + const result = WebgpuShaderProcessorWGSL.runCompute(device, source, {}, shader, reflectedGroupIndex); + return { result, shader }; +} + +describe('WebgpuShaderProcessorWGSL - compute reflection', function () { + + it('reflects a read_write storage buffer (regex fix)', function () { + const src = ` + var outBuf: array; + @compute @workgroup_size(1) fn main() { outBuf[0] = 1u; } + `; + const { result, shader } = run(src); + expect(shader.failed).to.equal(false); + + const fmt = result.computeBindGroupFormat; + expect(fmt.storageBufferFormats).to.have.lengthOf(1); + const sb = fmt.storageBufferFormats[0]; + expect(sb.name).to.equal('outBuf'); + expect(sb.readOnly).to.equal(false); + expect(sb.visibility).to.equal(SHADERSTAGE_COMPUTE); + expect(sb.slot).to.equal(0); + expect(result.computeUniformBufferFormat).to.equal(null); + + expect(result.cshader).to.contain('@group(1) @binding(0) var outBuf : array;'); + // original simplified declaration removed + expect(result.cshader).to.not.match(/^\s*var outBuf: array;/m); + }); + + it('reflects a read-only storage buffer', function () { + const src = ` + var inBuf: array; + @compute @workgroup_size(1) fn main() { let x = inBuf[0]; } + `; + const { result } = run(src); + const sb = result.computeBindGroupFormat.storageBufferFormats[0]; + expect(sb.readOnly).to.equal(true); + expect(result.cshader).to.contain('@group(1) @binding(0) var inBuf : array;'); + }); + + it('reflects a texture and its sampler (two slots)', function () { + const src = ` + var srcTex: texture_2d; + var srcTexSampler: sampler; + @compute @workgroup_size(1) fn main() { } + `; + const { result } = run(src); + const fmt = result.computeBindGroupFormat; + expect(fmt.textureFormats).to.have.lengthOf(1); + const tf = fmt.textureFormats[0]; + expect(tf.name).to.equal('srcTex'); + expect(tf.hasSampler).to.equal(true); + expect(tf.sampleType).to.equal(SAMPLETYPE_FLOAT); + expect(tf.slot).to.equal(0); + + expect(result.cshader).to.contain('@group(1) @binding(0) var srcTex: texture_2d;'); + expect(result.cshader).to.contain('@group(1) @binding(1) var srcTexSampler: sampler;'); + }); + + it('detects unfilterable-float (uff) textures', function () { + const src = ` + var dataTex: texture_2d; + @compute @workgroup_size(1) fn main() { } + `; + const { result } = run(src); + const tf = result.computeBindGroupFormat.textureFormats[0]; + expect(tf.sampleType).to.equal(SAMPLETYPE_UNFILTERABLE_FLOAT); + }); + + it('collapses loose uniforms into a single uniform buffer and renames access', function () { + const src = ` + uniform scale: f32; + uniform tint: vec4f; + @compute @workgroup_size(1) fn main() { let s = uniform.scale * uniform.tint.x; } + `; + const { result } = run(src); + + // one uniform buffer binding, no texture/storage + const ubf = result.computeUniformBufferFormat; + expect(ubf).to.not.equal(null); + expect(ubf.get('scale')).to.not.equal(undefined); + expect(ubf.get('tint')).to.not.equal(undefined); + expect(result.computeBindGroupFormat.uniformBufferFormats).to.have.lengthOf(1); + + // generated struct + buffer declaration + expect(result.cshader).to.contain('struct struct_ub_compute'); + expect(result.cshader).to.contain('var ub_compute : struct_ub_compute;'); + + // uniform.x references rewritten to ub_compute.x + expect(result.cshader).to.contain('ub_compute.scale'); + expect(result.cshader).to.contain('ub_compute.tint.x'); + expect(result.cshader).to.not.contain('uniform.scale'); + }); + + it('combines storage buffers, a texture, and uniforms in one reflected group', function () { + const src = ` + uniform scale: f32; + var srcTex: texture_2d; + var srcTexSampler: sampler; + var inBuf: array; + var outBuf: array; + @compute @workgroup_size(1) fn main() { } + `; + const { result } = run(src); + const fmt = result.computeBindGroupFormat; + expect(fmt.textureFormats).to.have.lengthOf(1); + expect(fmt.storageBufferFormats).to.have.lengthOf(2); + expect(fmt.uniformBufferFormats).to.have.lengthOf(1); + + // texture(0) + sampler(1) + 2 storage buffers(2,3) + uniform buffer(4) + expect(fmt.textureFormats[0].slot).to.equal(0); + const ubSlot = fmt.uniformBufferFormats[0].slot; + expect(ubSlot).to.equal(4); + expect(result.cshader).to.contain(`@group(1) @binding(${ubSlot}) var ub_compute`); + }); + + it('uses group 0 when no caller bind group format is supplied', function () { + const src = ` + var outBuf: array; + @compute @workgroup_size(1) fn main() { } + `; + const { result } = run(src, 0); + expect(result.cshader).to.contain('@group(0) @binding(0) var outBuf : array;'); + }); + + it('is a no-op when there is nothing to reflect (strictly additive)', function () { + const src = ` + @group(0) @binding(0) var outBuf: array; + @compute @workgroup_size(1) fn main() { outBuf[0] = 1u; } + `; + const { result } = run(src); + expect(result.computeBindGroupFormat).to.equal(null); + expect(result.computeUniformBufferFormat).to.equal(null); + expect(result.cshader).to.equal(src); + }); + + it('ignores workgroup, private and local variables', function () { + const src = ` + var wgScratch: array; + var counter: u32; + @compute @workgroup_size(1) fn main() { + var localVar: u32 = 0u; + let alias = wgScratch[0]; + } + `; + const { result } = run(src); + // nothing reflected + expect(result.computeBindGroupFormat).to.equal(null); + // declarations left untouched in the source + expect(result.cshader).to.contain('var wgScratch: array;'); + expect(result.cshader).to.contain('var counter: u32;'); + expect(result.cshader).to.contain('var localVar: u32 = 0u;'); + }); + + it('reflects a write-only storage texture', function () { + const src = ` + var outTex: texture_storage_2d; + @compute @workgroup_size(1) fn main() { } + `; + const { result } = run(src); + const fmt = result.computeBindGroupFormat; + expect(fmt.storageTextureFormats).to.have.lengthOf(1); + const st = fmt.storageTextureFormats[0]; + expect(st.name).to.equal('outTex'); + expect(st.write).to.equal(true); + expect(st.read).to.equal(false); + expect(st.slot).to.equal(0); + expect(result.cshader).to.contain('@group(1) @binding(0) var outTex: texture_storage_2d;'); + }); + + it('round-trips a read_write storage texture', function () { + const src = ` + var rwTex: texture_storage_2d; + @compute @workgroup_size(1) fn main() { } + `; + const { result } = run(src); + const st = result.computeBindGroupFormat.storageTextureFormats[0]; + expect(st.write).to.equal(true); + expect(st.read).to.equal(true); + expect(result.cshader).to.contain('@group(1) @binding(0) var rwTex: texture_storage_2d;'); + }); + + it('reflects a sampled texture + sampler + storage texture (edge-detect layout, group 0)', function () { + const src = ` + var inputTexture: texture_2d; + var inputTexture_sampler: sampler; + var outputTexture: texture_storage_2d; + @compute @workgroup_size(8, 8, 1) fn main() { } + `; + const { result } = run(src, 0); + const fmt = result.computeBindGroupFormat; + expect(fmt.textureFormats).to.have.lengthOf(1); + expect(fmt.storageTextureFormats).to.have.lengthOf(1); + + // inputTexture(0) + sampler(1) + outputTexture(2) + expect(fmt.textureFormats[0].slot).to.equal(0); + expect(fmt.storageTextureFormats[0].slot).to.equal(2); + + expect(result.cshader).to.contain('@group(0) @binding(0) var inputTexture: texture_2d;'); + expect(result.cshader).to.contain('@group(0) @binding(1) var inputTexture_sampler: sampler;'); + expect(result.cshader).to.contain('@group(0) @binding(2) var outputTexture: texture_storage_2d;'); + }); + + it('reflects the particles-simulation pattern (loose uniforms + rw/ro storage, group 0)', function () { + const src = ` + struct Sphere { center: vec3, radius: f32 } + uniform count: u32; + uniform dt: f32; + uniform sphereCount: u32; + var particles: array; + var spheres: array; + @compute @workgroup_size(64) fn main() { let n = uniform.count; } + `; + const { result } = run(src, 0); + const fmt = result.computeBindGroupFormat; + + // two storage buffers + one generated uniform buffer + expect(fmt.storageBufferFormats.map(f => f.name)).to.have.members(['particles', 'spheres']); + expect(fmt.uniformBufferFormats).to.have.lengthOf(1); + + const particlesFmt = fmt.storageBufferFormats.find(f => f.name === 'particles'); + const spheresFmt = fmt.storageBufferFormats.find(f => f.name === 'spheres'); + expect(particlesFmt.readOnly).to.equal(false); + expect(spheresFmt.readOnly).to.equal(true); + + // particles(0), spheres(1), ub_compute(2) + expect(result.cshader).to.contain('@group(0) @binding(0) var particles : array;'); + expect(result.cshader).to.contain('@group(0) @binding(1) var spheres : array;'); + expect(result.cshader).to.contain('@group(0) @binding(2) var ub_compute'); + expect(result.cshader).to.contain('ub_compute.count'); + }); + + it('leaves explicitly-bound resources untouched while reflecting simplified ones', function () { + const src = ` + @group(0) @binding(0) var manualOut: array; + uniform scale: f32; + var reflectedIn: array; + @compute @workgroup_size(1) fn main() { manualOut[0] = u32(uniform.scale); } + `; + const { result } = run(src); + + // explicit binding stays exactly as written, in group 0 + expect(result.cshader).to.contain('@group(0) @binding(0) var manualOut: array;'); + + // reflected resources go to group 1 + const fmt = result.computeBindGroupFormat; + expect(fmt.storageBufferFormats.map(f => f.name)).to.include('reflectedIn'); + expect(fmt.storageBufferFormats.map(f => f.name)).to.not.include('manualOut'); + expect(result.cshader).to.contain('@group(1) @binding(0) var reflectedIn : array;'); + expect(result.cshader).to.contain('ub_compute.scale'); + }); + +}); From 2e98f17e03cf866de3623c6e33548a77300c7978 Mon Sep 17 00:00:00 2001 From: Martin Valigursky Date: Tue, 2 Jun 2026 09:48:55 +0100 Subject: [PATCH 2/3] fix(webgpu): avoid GPUComputePipeline type in generated d.ts The @returns {GPUComputePipeline} JSDoc emitted an explicit return type into the .d.ts, which fails the standalone type check (no WebGPU lib types in scope). --- src/platform/graphics/webgpu/webgpu-compute-pipeline.js | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/src/platform/graphics/webgpu/webgpu-compute-pipeline.js b/src/platform/graphics/webgpu/webgpu-compute-pipeline.js index 6f5db71d378..57ecf4de5a1 100644 --- a/src/platform/graphics/webgpu/webgpu-compute-pipeline.js +++ b/src/platform/graphics/webgpu/webgpu-compute-pipeline.js @@ -43,7 +43,7 @@ class WebgpuComputePipeline extends WebgpuPipeline { * @param {import('../shader.js').Shader} shader - The compute shader. * @param {import('../bind-group-format.js').BindGroupFormat[]} bindGroupFormats - The bind group * formats, in bind group index order (dense, no gaps). - * @returns {GPUComputePipeline} - The compute pipeline. + * @returns {object} - The compute pipeline (GPUComputePipeline). */ get(shader, bindGroupFormats) { From d98d94776ac9b32db8119698582e112c8ec5b120 Mon Sep 17 00:00:00 2001 From: Martin Valigursky Date: Tue, 2 Jun 2026 10:49:16 +0100 Subject: [PATCH 3/3] refactor(webgpu): address compute reflection review comments - gpuFormatToPixelFormat: use last-wins so the canonical PIXELFORMAT (e.g. RGBA8 for 'rgba8unorm') is chosen, avoiding a needless bind group layout / pipeline cache split (the key uses the numeric PIXELFORMAT). - Assert computeUniformBufferFormats requires a computeBindGroupFormat, instead of silently dropping the uniform buffers. - STORAGE_BUFFER_REGEX: drop the invalid 'write' token (storage buffers only support read / read_write in WGSL). - Sync runCompute and buildResourceFormats JSDoc to mention storage textures. --- .../graphics/webgpu/webgpu-compute.js | 4 +++ .../webgpu/webgpu-shader-processor-wgsl.js | 25 +++++++++++-------- 2 files changed, 18 insertions(+), 11 deletions(-) diff --git a/src/platform/graphics/webgpu/webgpu-compute.js b/src/platform/graphics/webgpu/webgpu-compute.js index 28b68b44daa..1571b165f7f 100644 --- a/src/platform/graphics/webgpu/webgpu-compute.js +++ b/src/platform/graphics/webgpu/webgpu-compute.js @@ -37,6 +37,10 @@ class WebgpuCompute { computeReflectedGroupIndex } = shader.impl; + // caller uniform buffers are bound into the caller bind group, so the format is required + Debug.assert(!computeUniformBufferFormats || computeBindGroupFormat, + 'Compute shader specifies computeUniformBufferFormats but no computeBindGroupFormat to bind them into', shader); + // ordered, gapless array of bind group formats (array index === bind group index) const formats = []; diff --git a/src/platform/graphics/webgpu/webgpu-shader-processor-wgsl.js b/src/platform/graphics/webgpu/webgpu-shader-processor-wgsl.js index 2f6c6a43e5e..592bc98bec6 100644 --- a/src/platform/graphics/webgpu/webgpu-shader-processor-wgsl.js +++ b/src/platform/graphics/webgpu/webgpu-shader-processor-wgsl.js @@ -176,12 +176,13 @@ const getTextureDeclarationType = (viewDimension, sampleType) => { // reverse of gpuTextureFormats: WGSL/GPU storage format string -> PIXELFORMAT. Built once, used to // reflect storage texture declarations. Several PIXELFORMATs can map to the same string (e.g. RGB8 -// and RGBA8 both -> 'rgba8unorm'); first-wins is fine as the chosen PIXELFORMAT only needs to -// round-trip back to the same string via gpuTextureFormats[...], which is what the bind group -// layout (and WebGPU validation) uses. +// and RGBA8 both -> 'rgba8unorm'); last-wins so the canonical/most-common format wins (RGBA8 over +// RGB8 with the current table). This matches what callers pass to a hand-authored +// BindStorageTextureFormat, avoiding a needless split of the bind group layout / pipeline cache +// (whose key uses the numeric PIXELFORMAT, even though the GPUTextureFormat string is identical). const gpuFormatToPixelFormat = new Map(); gpuTextureFormats.forEach((str, pixelFormat) => { - if (str && !gpuFormatToPixelFormat.has(str)) { + if (str) { gpuFormatToPixelFormat.set(str, pixelFormat); } }); @@ -265,8 +266,9 @@ class UniformLine { const TEXTURE_REGEX = /^\s*var\s+(\w+)\s*:\s*(texture_\w+)(?:<(\w+)>)?;\s*$/; // eslint-disable-next-line const STORAGE_TEXTURE_REGEX = /^\s*var\s+([\w\d_]+)\s*:\s*(texture_storage_2d|texture_storage_2d_array)<([\w\d_]+),\s*(\w+)>\s*;\s*$/; +// storage buffers only support 'read' and 'read_write' access in WGSL (no write-only form) // eslint-disable-next-line -const STORAGE_BUFFER_REGEX = /^\s*var\s*\s*([\w\d_]+)\s*:\s*(.*)\s*;\s*$/; +const STORAGE_BUFFER_REGEX = /^\s*var\s*\s*([\w\d_]+)\s*:\s*(.*)\s*;\s*$/; // eslint-disable-next-line const EXTERNAL_TEXTURE_REGEX = /^\s*var\s+([\w\d_]+)\s*:\s*texture_external;\s*$/; // eslint-disable-next-line @@ -460,7 +462,8 @@ class WebgpuShaderProcessorWGSL { /** * Process a compute shader: reflect its simplified-syntax declarations (loose `uniform`s, - * textures/samplers, storage buffers) into a single bind group at `reflectedGroupIndex`, leaving + * textures/samplers, storage buffers, storage textures) into a single bind group at + * `reflectedGroupIndex`, leaving * any explicitly-bound (`@group/@binding`) declarations untouched. The loose uniforms are * collapsed into one generated uniform buffer (`ub_compute`) placed inside that same group. * @@ -699,15 +702,15 @@ class WebgpuShaderProcessorWGSL { } /** - * Converts parsed resource lines (textures, samplers, storage buffers) into an array of bind - * formats. Shared by the vertex/fragment path ({@link processResources}) and the compute path - * ({@link runCompute}); only the shader-stage visibility differs. + * Converts parsed resource lines (textures, samplers, storage buffers, storage textures) into an + * array of bind formats. Shared by the vertex/fragment path ({@link processResources}) and the + * compute path ({@link runCompute}); only the shader-stage visibility differs. * * @param {Array} resources - The parsed resource lines. * @param {number} visibility - Shader stage visibility bit-flags for the created formats. * @param {Shader} shader - The shader (for error reporting). - * @returns {Array} - The bind formats, in declaration - * order (a texture with a sampler consumes the following sampler line). + * @returns {Array} - The bind + * formats, in declaration order (a texture with a sampler consumes the following sampler line). * @private */ static buildResourceFormats(resources, visibility, shader) {