From b1fe7627038dd8460fb2fefa6d32dd5a8ff27d61 Mon Sep 17 00:00:00 2001
From: sunag <jcdeconto@igalia.com>
Date: Sat, 16 Apr 2022 14:15:53 -0300
Subject: [PATCH 01/11] add ComputeNode

---
 examples/jsm/nodes/Nodes.js             | 10 +++-
 examples/jsm/nodes/gpgpu/ComputeNode.js | 62 +++++++++++++++++++++++++
 2 files changed, 71 insertions(+), 1 deletion(-)
 create mode 100644 examples/jsm/nodes/gpgpu/ComputeNode.js

diff --git a/examples/jsm/nodes/Nodes.js b/examples/jsm/nodes/Nodes.js
index 4421c4dcabae9b..38c9a68eda5658 100644
--- a/examples/jsm/nodes/Nodes.js
+++ b/examples/jsm/nodes/Nodes.js
@@ -44,6 +44,9 @@ import SkinningNode from './accessors/SkinningNode.js';
 import TextureNode from './accessors/TextureNode.js';
 import UVNode from './accessors/UVNode.js';
 
+// gpgpu
+import ComputeNode from './gpgpu/ComputeNode.js';
+
 // display
 import ColorSpaceNode from './display/ColorSpaceNode.js';
 import NormalMapNode from './display/NormalMapNode.js';
@@ -118,6 +121,9 @@ const nodeLib = {
 	VarNode,
 	VaryNode,
 
+	// compute
+	ComputeNode,
+
 	// accessors
 	BufferNode,
 	CameraNode,
@@ -210,6 +216,9 @@ export {
 	VarNode,
 	VaryNode,
 
+	// compute
+	ComputeNode,
+
 	// accessors
 	BufferNode,
 	CameraNode,
@@ -265,5 +274,4 @@ export {
 	NodeLoader,
 	NodeObjectLoader,
 	NodeMaterialLoader
-
 };
diff --git a/examples/jsm/nodes/gpgpu/ComputeNode.js b/examples/jsm/nodes/gpgpu/ComputeNode.js
new file mode 100644
index 00000000000000..4718613e685807
--- /dev/null
+++ b/examples/jsm/nodes/gpgpu/ComputeNode.js
@@ -0,0 +1,62 @@
+import Node from '../core/Node.js';
+import { assign, element, instanceIndex } from '../shadernode/ShaderNodeElements.js';
+import { NodeUpdateType } from '../core/constants.js';
+
+class ComputeNode extends Node {
+
+	constructor( dispatchCount, workgroupSize = [ 64 ] ) {
+
+		super( 'void' );
+
+		this.updateType = NodeUpdateType.Object;
+
+		this.dispatchCount = dispatchCount;
+		this.workgroupSize = workgroupSize;
+
+		this.assigns = [];
+
+	}
+
+	getMainStorageBufferNode() {
+
+		const assigns = this.assigns;
+
+		return assigns[ assigns.length - 1 ].storageBufferNode;
+
+	}
+
+	assign( storageBufferNode, sourceNode ) {
+
+		this.assigns.push( { storageBufferNode, sourceNode } );
+
+		return this;
+
+	}
+
+	update( { renderer } ) {
+
+		renderer.compute( this );
+
+	}
+
+	generate( builder ) {
+
+		const { renderer, shaderStage } = builder;
+
+		if ( shaderStage === 'compute' ) {
+
+			for ( const { storageBufferNode, sourceNode } of this.assigns ) {
+
+				assign( element( storageBufferNode, instanceIndex ), sourceNode ).build( builder );
+
+			}
+
+		}
+
+	}
+
+}
+
+ComputeNode.prototype.isComputeNode = true;
+
+export default ComputeNode;

From 0253faf62c553ff15ad49a9216ebfb538be0deda Mon Sep 17 00:00:00 2001
From: sunag <jcdeconto@igalia.com>
Date: Sat, 16 Apr 2022 14:18:00 -0300
Subject: [PATCH 02/11] fix float function parameters, support to void function

---
 examples/jsm/nodes/parsers/WGSLNodeFunction.js | 16 ++++++++++++----
 1 file changed, 12 insertions(+), 4 deletions(-)

diff --git a/examples/jsm/nodes/parsers/WGSLNodeFunction.js b/examples/jsm/nodes/parsers/WGSLNodeFunction.js
index d442b13812218a..0119477e37ef70 100644
--- a/examples/jsm/nodes/parsers/WGSLNodeFunction.js
+++ b/examples/jsm/nodes/parsers/WGSLNodeFunction.js
@@ -1,9 +1,13 @@
 import NodeFunction from '../core/NodeFunction.js';
 import NodeFunctionInput from '../core/NodeFunctionInput.js';
 
-const declarationRegexp = /^fn\s*([a-z_0-9]+)?\s*\(([\s\S]*?)\)\s*\-\>\s*([a-z_0-9]+)?/i;
+const declarationRegexp = /^[fn]*\s*([a-z_0-9]+)?\s*\(([\s\S]*?)\)\s*[\-\>]*\s*([a-z_0-9]+)?/i;
 const propertiesRegexp = /[a-z_0-9]+/ig;
 
+const wgslTypeLib = {
+	f32: 'float'
+};
+
 const parse = ( source ) => {
 
 	source = source.trim();
@@ -36,7 +40,9 @@ const parse = ( source ) => {
 			// default
 
 			const name = propsMatches[ i ++ ][ 0 ];
-			const type = propsMatches[ i ++ ][ 0 ];
+			let type = propsMatches[ i ++ ][ 0 ];
+
+			type = wgslTypeLib[ type ] || type;
 
 			// precision
 
@@ -54,7 +60,7 @@ const parse = ( source ) => {
 		const blockCode = source.substring( declaration[ 0 ].length );
 
 		const name = declaration[ 1 ] !== undefined ? declaration[ 1 ] : '';
-		const type = declaration[ 3 ];
+		const type = declaration[ 3 ] || 'void';
 
 		return {
 			type,
@@ -87,7 +93,9 @@ class WGSLNodeFunction extends NodeFunction {
 
 	getCode( name = this.name ) {
 
-		return `fn ${ name } ( ${ this.inputsCode.trim() } ) -> ${ this.type }` + this.blockCode;
+		const type = this.type !== 'void' ? '-> ' + this.type : '';
+
+		return `fn ${ name } ( ${ this.inputsCode.trim() } ) ${ type }` + this.blockCode;
 
 	}
 

From 83c85c6e9c72f034c9302c15a801db976690ca6d Mon Sep 17 00:00:00 2001
From: sunag <jcdeconto@igalia.com>
Date: Sat, 16 Apr 2022 14:18:22 -0300
Subject: [PATCH 03/11] add StorageBufferNode

---
 .../jsm/nodes/accessors/StorageBufferNode.js  | 21 +++++++++++++++++++
 1 file changed, 21 insertions(+)
 create mode 100644 examples/jsm/nodes/accessors/StorageBufferNode.js

diff --git a/examples/jsm/nodes/accessors/StorageBufferNode.js b/examples/jsm/nodes/accessors/StorageBufferNode.js
new file mode 100644
index 00000000000000..e000bdb3c5ea5c
--- /dev/null
+++ b/examples/jsm/nodes/accessors/StorageBufferNode.js
@@ -0,0 +1,21 @@
+import BufferNode from './BufferNode.js';
+
+class StorageBufferNode extends BufferNode {
+
+	constructor( value, bufferType, bufferCount = 0 ) {
+
+		super( value, bufferType, bufferCount );
+
+	}
+
+	getInputType( /*builder*/ ) {
+
+		return 'storageBuffer';
+
+	}
+
+}
+
+StorageBufferNode.prototype.isStorageBufferNode = true;
+
+export default StorageBufferNode;

From bf2db826b5bcd88d6c196729162f752c46228bc6 Mon Sep 17 00:00:00 2001
From: sunag <jcdeconto@igalia.com>
Date: Sat, 16 Apr 2022 14:19:33 -0300
Subject: [PATCH 04/11] add compute(), timer(), compute(), storage() and func()
 elements

---
 examples/jsm/nodes/shadernode/ShaderNodeElements.js | 12 ++++++++++++
 1 file changed, 12 insertions(+)

diff --git a/examples/jsm/nodes/shadernode/ShaderNodeElements.js b/examples/jsm/nodes/shadernode/ShaderNodeElements.js
index cb11cd525a6c7a..164a48c811e29f 100644
--- a/examples/jsm/nodes/shadernode/ShaderNodeElements.js
+++ b/examples/jsm/nodes/shadernode/ShaderNodeElements.js
@@ -6,9 +6,11 @@ import UniformNode from '../core/UniformNode.js';
 import BypassNode from '../core/BypassNode.js';
 import InstanceIndexNode from '../core/InstanceIndexNode.js';
 import ContextNode from '../core/ContextNode.js';
+import FunctionNode from '../core/FunctionNode.js';
 
 // accessor nodes
 import BufferNode from '../accessors/BufferNode.js';
+import StorageBufferNode from '../accessors/StorageBufferNode.js';
 import CameraNode from '../accessors/CameraNode.js';
 import MaterialNode from '../accessors/MaterialNode.js';
 import ModelNode from '../accessors/ModelNode.js';
@@ -20,6 +22,9 @@ import TextureNode from '../accessors/TextureNode.js';
 import UVNode from '../accessors/UVNode.js';
 import InstanceNode from '../accessors/InstanceNode.js';
 
+// gpgpu
+import ComputeNode from '../gpgpu/ComputeNode.js';
+
 // math nodes
 import OperatorNode from '../math/OperatorNode.js';
 import CondNode from '../math/CondNode.js';
@@ -29,6 +34,7 @@ import MathNode from '../math/MathNode.js';
 import ArrayElementNode from '../utils/ArrayElementNode.js';
 import ConvertNode from '../utils/ConvertNode.js';
 import JoinNode from '../utils/JoinNode.js';
+import TimerNode from '../utils/TimerNode.js';
 
 // other nodes
 import ColorSpaceNode from '../display/ColorSpaceNode.js';
@@ -109,9 +115,15 @@ export const join = ( ...params ) => nodeObject( new JoinNode( nodeArray( params
 export const uv = ( ...params ) => nodeObject( new UVNode( ...params ) );
 export const attribute = ( ...params ) => nodeObject( new AttributeNode( ...params ) );
 export const buffer = ( ...params ) => nodeObject( new BufferNode( ...params ) );
+export const storage = ( ...params ) => nodeObject( new StorageBufferNode( ...params ) );
 export const texture = ( ...params ) => nodeObject( new TextureNode( ...params ) );
 export const sampler = ( texture ) => nodeObject( new ConvertNode( texture.isNode === true ? texture : new TextureNode( texture ), 'sampler' ) );
 
+export const timer = ( ...params ) => nodeObject( new TimerNode( ...params ) );
+
+export const compute = ( ...params ) => nodeObject( new ComputeNode( ...params ) );
+export const func = ( ...params ) => nodeObject( new FunctionNode( ...params ) );
+
 export const cond = nodeProxy( CondNode );
 
 export const add = nodeProxy( OperatorNode, '+' );

From c0f702204ba0ed88a75e609938938fcd053fbdef Mon Sep 17 00:00:00 2001
From: sunag <jcdeconto@igalia.com>
Date: Sat, 16 Apr 2022 14:20:55 -0300
Subject: [PATCH 05/11] WebGPU: add StorageBufferNode and ComputeNode support
 and cleanup

---
 examples/jsm/nodes/core/NodeBuilder.js        |  63 ++++++----
 .../jsm/renderers/webgpu/WebGPUBindings.js    |  13 +-
 examples/jsm/renderers/webgpu/WebGPUBuffer.js |  43 +++++++
 .../webgpu/WebGPUComputePipelines.js          |  18 ++-
 .../renderers/webgpu/WebGPURenderPipelines.js |   3 +-
 .../jsm/renderers/webgpu/WebGPURenderer.js    |  18 +--
 .../renderers/webgpu/WebGPUStorageBuffer.js   |  11 +-
 .../renderers/webgpu/WebGPUUniformBuffer.js   |  35 +-----
 .../webgpu/nodes/WebGPUNodeBuilder.js         | 111 ++++++++++++++----
 .../webgpu/nodes/WebGPUNodeUniformsGroup.js   |  20 ----
 10 files changed, 210 insertions(+), 125 deletions(-)
 create mode 100644 examples/jsm/renderers/webgpu/WebGPUBuffer.js
 delete mode 100644 examples/jsm/renderers/webgpu/nodes/WebGPUNodeUniformsGroup.js

diff --git a/examples/jsm/nodes/core/NodeBuilder.js b/examples/jsm/nodes/core/NodeBuilder.js
index 2fe31645b8a531..02fe433b42cfdf 100644
--- a/examples/jsm/nodes/core/NodeBuilder.js
+++ b/examples/jsm/nodes/core/NodeBuilder.js
@@ -24,7 +24,7 @@ class NodeBuilder {
 	constructor( object, renderer, parser ) {
 
 		this.object = object;
-		this.material = object.material;
+		this.material = object.material || null;
 		this.renderer = renderer;
 		this.parser = parser;
 
@@ -34,14 +34,15 @@ class NodeBuilder {
 
 		this.vertexShader = null;
 		this.fragmentShader = null;
+		this.computeShader = null;
 
-		this.flowNodes = { vertex: [], fragment: [] };
-		this.flowCode = { vertex: '', fragment: '' };
-		this.uniforms = { vertex: [], fragment: [], index: 0 };
-		this.codes = { vertex: [], fragment: [] };
+		this.flowNodes = { vertex: [], fragment: [], compute: [] };
+		this.flowCode = { vertex: '', fragment: '', compute: [] };
+		this.uniforms = { vertex: [], fragment: [], compute: [], index: 0 };
+		this.codes = { vertex: [], fragment: [], compute: [] };
 		this.attributes = [];
 		this.varys = [];
-		this.vars = { vertex: [], fragment: [] };
+		this.vars = { vertex: [], fragment: [], compute: [] };
 		this.flow = { code: '' };
 		this.stack = [];
 
@@ -369,7 +370,7 @@ class NodeBuilder {
 
 		if ( nodeData === undefined ) {
 
-			nodeData = { vertex: {}, fragment: {} };
+			nodeData = { vertex: {}, fragment: {}, compute: {} };
 
 			this.nodesData.set( node, nodeData );
 
@@ -592,7 +593,7 @@ class NodeBuilder {
 
 	getHash() {
 
-		return this.vertexShader + this.fragmentShader;
+		return this.vertexShader + this.fragmentShader + this.computeShader;
 
 	}
 
@@ -616,33 +617,51 @@ class NodeBuilder {
 
 	build() {
 
-		// stage 1: analyze nodes to possible optimization and validation
+		if ( this.material !== null ) {
 
-		for ( const shaderStage of shaderStages ) {
+			// stage 1: analyze nodes to possible optimization and validation
 
-			this.setShaderStage( shaderStage );
+			for ( const shaderStage of shaderStages ) {
 
-			const flowNodes = this.flowNodes[ shaderStage ];
+				this.setShaderStage( shaderStage );
 
-			for ( const node of flowNodes ) {
+				const flowNodes = this.flowNodes[ shaderStage ];
 
-				node.analyze( this );
+				for ( const node of flowNodes ) {
+
+					node.analyze( this );
+
+				}
 
 			}
 
-		}
+			// stage 2: pre-build vertex code used in fragment shader
 
-		// stage 2: pre-build vertex code used in fragment shader
+			if ( this.context.vertex && this.context.vertex.isNode ) {
 
-		if ( this.context.vertex && this.context.vertex.isNode ) {
+				this.flowNodeFromShaderStage( 'vertex', this.context.vertex );
 
-			this.flowNodeFromShaderStage( 'vertex', this.context.vertex );
+			}
 
-		}
+			// stage 3: generate shader
+
+			for ( const shaderStage of shaderStages ) {
 
-		// stage 3: generate shader
+				this.setShaderStage( shaderStage );
 
-		for ( const shaderStage of shaderStages ) {
+				const flowNodes = this.flowNodes[ shaderStage ];
+
+				for ( const node of flowNodes ) {
+
+					this.flowNode( node, shaderStage );
+
+				}
+
+			}
+
+		} else {
+
+			const shaderStage = 'compute';
 
 			this.setShaderStage( shaderStage );
 
@@ -650,6 +669,8 @@ class NodeBuilder {
 
 			for ( const node of flowNodes ) {
 
+				node.analyze( this );
+
 				this.flowNode( node, shaderStage );
 
 			}
diff --git a/examples/jsm/renderers/webgpu/WebGPUBindings.js b/examples/jsm/renderers/webgpu/WebGPUBindings.js
index 35605d314067a3..9e5810fbd9643b 100644
--- a/examples/jsm/renderers/webgpu/WebGPUBindings.js
+++ b/examples/jsm/renderers/webgpu/WebGPUBindings.js
@@ -30,9 +30,9 @@ class WebGPUBindings {
 
 			// setup (static) binding layout and (dynamic) binding group
 
-			const renderPipeline = this.renderPipelines.get( object );
+			const pipeline = object.isNode ? this.computePipelines.get( object ) : this.renderPipelines.get( object ).pipeline;
 
-			const bindLayout = renderPipeline.pipeline.getBindGroupLayout( 0 );
+			const bindLayout = pipeline.getBindGroupLayout( 0 );
 			const bindGroup = this._createBindGroup( bindings, bindLayout );
 
 			data = {
@@ -108,12 +108,12 @@ class WebGPUBindings {
 			if ( binding.isUniformBuffer ) {
 
 				const buffer = binding.getBuffer();
-				const bufferGPU = binding.bufferGPU;
-
 				const needsBufferWrite = binding.update();
 
 				if ( needsBufferWrite === true ) {
 
+					const bufferGPU = binding.bufferGPU;
+
 					this.device.queue.writeBuffer( bufferGPU, 0, buffer, 0 );
 
 				}
@@ -121,6 +121,7 @@ class WebGPUBindings {
 			} else if ( binding.isStorageBuffer ) {
 
 				const attribute = binding.attribute;
+
 				this.attributes.update( attribute, false, binding.usage );
 
 			} else if ( binding.isSampler ) {
@@ -243,8 +244,8 @@ class WebGPUBindings {
 		}
 
 		return this.device.createBindGroup( {
-			layout: layout,
-			entries: entries
+			layout,
+			entries
 		} );
 
 	}
diff --git a/examples/jsm/renderers/webgpu/WebGPUBuffer.js b/examples/jsm/renderers/webgpu/WebGPUBuffer.js
new file mode 100644
index 00000000000000..072dbadb6c23d3
--- /dev/null
+++ b/examples/jsm/renderers/webgpu/WebGPUBuffer.js
@@ -0,0 +1,43 @@
+import WebGPUBinding from './WebGPUBinding.js';
+import { getFloatLength } from './WebGPUBufferUtils.js';
+
+class WebGPUBuffer extends WebGPUBinding {
+
+	constructor( name, type, buffer = null ) {
+
+		super( name );
+
+		this.bytesPerElement = Float32Array.BYTES_PER_ELEMENT;
+		this.type = type;
+		this.visibility = GPUShaderStage.VERTEX;
+
+		this.usage = GPUBufferUsage.COPY_DST;
+
+		this.buffer = buffer;
+		this.bufferGPU = null; // set by the renderer
+
+	}
+
+	getByteLength() {
+
+		return getFloatLength( this.buffer.byteLength );
+
+	}
+
+	getBuffer() {
+
+		return this.buffer;
+
+	}
+
+	update() {
+
+		return true;
+
+	}
+
+}
+
+WebGPUBuffer.prototype.isBuffer = true;
+
+export default WebGPUBuffer;
diff --git a/examples/jsm/renderers/webgpu/WebGPUComputePipelines.js b/examples/jsm/renderers/webgpu/WebGPUComputePipelines.js
index bd604aac006137..8c9f952e4367c8 100644
--- a/examples/jsm/renderers/webgpu/WebGPUComputePipelines.js
+++ b/examples/jsm/renderers/webgpu/WebGPUComputePipelines.js
@@ -2,9 +2,10 @@ import WebGPUProgrammableStage from './WebGPUProgrammableStage.js';
 
 class WebGPUComputePipelines {
 
-	constructor( device ) {
+	constructor( device, nodes ) {
 
 		this.device = device;
+		this.nodes = nodes;
 
 		this.pipelines = new WeakMap();
 		this.stages = {
@@ -13,9 +14,9 @@ class WebGPUComputePipelines {
 
 	}
 
-	get( param ) {
+	get( computeNode ) {
 
-		let pipeline = this.pipelines.get( param );
+		let pipeline = this.pipelines.get( computeNode );
 
 		// @TODO: Reuse compute pipeline if possible, introduce WebGPUComputePipeline
 
@@ -23,8 +24,13 @@ class WebGPUComputePipelines {
 
 			const device = this.device;
 
+			// get shader
+
+			const nodeBuilder = this.nodes.get( computeNode );
+			const computeShader = nodeBuilder.computeShader;
+
 			const shader = {
-				computeShader: param.shader
+				computeShader
 			};
 
 			// programmable stage
@@ -33,7 +39,7 @@ class WebGPUComputePipelines {
 
 			if ( stageCompute === undefined ) {
 
- 				stageCompute = new WebGPUProgrammableStage( device, shader.computeShader, 'compute' );
+ 				stageCompute = new WebGPUProgrammableStage( device, computeShader, 'compute' );
 
 				this.stages.compute.set( shader, stageCompute );
 
@@ -43,7 +49,7 @@ class WebGPUComputePipelines {
 				compute: stageCompute.stage
 			} );
 
-			this.pipelines.set( param, pipeline );
+			this.pipelines.set( computeNode, pipeline );
 
 		}
 
diff --git a/examples/jsm/renderers/webgpu/WebGPURenderPipelines.js b/examples/jsm/renderers/webgpu/WebGPURenderPipelines.js
index 1196a2996ef830..8e938c775844d9 100644
--- a/examples/jsm/renderers/webgpu/WebGPURenderPipelines.js
+++ b/examples/jsm/renderers/webgpu/WebGPURenderPipelines.js
@@ -24,7 +24,6 @@ class WebGPURenderPipelines {
 	get( object ) {
 
 		const device = this.device;
-		const material = object.material;
 
 		const cache = this._getCache( object );
 
@@ -32,6 +31,8 @@ class WebGPURenderPipelines {
 
 		if ( this._needsUpdate( object, cache ) ) {
 
+			const material = object.material;
+
 			// release previous cache
 
 			if ( cache.currentPipeline !== undefined ) {
diff --git a/examples/jsm/renderers/webgpu/WebGPURenderer.js b/examples/jsm/renderers/webgpu/WebGPURenderer.js
index eeeb2db2961b77..71abb7fba99c7e 100644
--- a/examples/jsm/renderers/webgpu/WebGPURenderer.js
+++ b/examples/jsm/renderers/webgpu/WebGPURenderer.js
@@ -186,7 +186,7 @@ class WebGPURenderer {
 		this._textures = new WebGPUTextures( device, this._properties, this._info );
 		this._objects = new WebGPUObjects( this._geometries, this._info );
 		this._nodes = new WebGPUNodes( this, this._properties );
-		this._computePipelines = new WebGPUComputePipelines( device );
+		this._computePipelines = new WebGPUComputePipelines( device, this._nodes );
 		this._renderPipelines = new WebGPURenderPipelines( this, device, parameters.sampleCount, this._nodes );
 		this._bindings = this._renderPipelines.bindings = new WebGPUBindings( device, this._info, this._properties, this._textures, this._renderPipelines, this._computePipelines, this._attributes, this._nodes );
 		this._renderLists = new WebGPURenderLists();
@@ -611,26 +611,30 @@ class WebGPURenderer {
 
 	}
 
-	compute( computeParams ) {
+	compute( ...computeNodes ) {
 
 		const device = this._device;
 		const cmdEncoder = device.createCommandEncoder( {} );
 		const passEncoder = cmdEncoder.beginComputePass();
 
-		for ( const param of computeParams ) {
+		for ( const computeNode of computeNodes ) {
 
 			// pipeline
 
-			const pipeline = this._computePipelines.get( param );
+			const pipeline = this._computePipelines.get( computeNode );
 			passEncoder.setPipeline( pipeline );
 
+			// node
+
+			//this._nodes.update( computeNode );
+
 			// bind group
 
-			const bindGroup = this._bindings.getForCompute( param ).group;
-			this._bindings.update( param );
+			const bindGroup = this._bindings.get( computeNode ).group;
+			this._bindings.update( computeNode );
 			passEncoder.setBindGroup( 0, bindGroup );
 
-			passEncoder.dispatch( param.num );
+			passEncoder.dispatch( computeNode.dispatchCount );
 
 		}
 
diff --git a/examples/jsm/renderers/webgpu/WebGPUStorageBuffer.js b/examples/jsm/renderers/webgpu/WebGPUStorageBuffer.js
index 95eb34fbb5aec4..07868a9e7a1bdd 100644
--- a/examples/jsm/renderers/webgpu/WebGPUStorageBuffer.js
+++ b/examples/jsm/renderers/webgpu/WebGPUStorageBuffer.js
@@ -1,18 +1,15 @@
-import WebGPUBinding from './WebGPUBinding.js';
+import WebGPUBuffer from './WebGPUBuffer.js';
 import { GPUBindingType } from './constants.js';
 
-class WebGPUStorageBuffer extends WebGPUBinding {
+class WebGPUStorageBuffer extends WebGPUBuffer {
 
 	constructor( name, attribute ) {
 
-		super( name );
+		super( name, GPUBindingType.StorageBuffer, attribute.array );
 
-		this.type = GPUBindingType.StorageBuffer;
-
-		this.usage = GPUBufferUsage.VERTEX | GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST;
+		this.usage |= GPUBufferUsage.VERTEX | GPUBufferUsage.STORAGE;
 
 		this.attribute = attribute;
-		this.bufferGPU = null; // set by the renderer
 
 	}
 
diff --git a/examples/jsm/renderers/webgpu/WebGPUUniformBuffer.js b/examples/jsm/renderers/webgpu/WebGPUUniformBuffer.js
index 1d47da27e9629c..cc03925e584d20 100644
--- a/examples/jsm/renderers/webgpu/WebGPUUniformBuffer.js
+++ b/examples/jsm/renderers/webgpu/WebGPUUniformBuffer.js
@@ -1,40 +1,13 @@
-import WebGPUBinding from './WebGPUBinding.js';
-import { getFloatLength } from './WebGPUBufferUtils.js';
-
+import WebGPUBuffer from './WebGPUBuffer.js';
 import { GPUBindingType } from './constants.js';
 
-class WebGPUUniformBuffer extends WebGPUBinding {
+class WebGPUUniformBuffer extends WebGPUBuffer {
 
 	constructor( name, buffer = null ) {
 
-		super( name );
-
-		this.bytesPerElement = Float32Array.BYTES_PER_ELEMENT;
-		this.type = GPUBindingType.UniformBuffer;
-		this.visibility = GPUShaderStage.VERTEX;
-
-		this.usage = GPUBufferUsage.UNIFORM | GPUBufferUsage.COPY_DST;
-
-		this.buffer = buffer;
-		this.bufferGPU = null; // set by the renderer
-
-	}
-
-	getByteLength() {
-
-		return getFloatLength( this.buffer.byteLength );
-
-	}
-
-	getBuffer() {
-
-		return this.buffer;
-
-	}
-
-	update() {
+		super( name, GPUBindingType.UniformBuffer, buffer );
 
-		return true;
+		this.usage |= GPUBufferUsage.UNIFORM;
 
 	}
 
diff --git a/examples/jsm/renderers/webgpu/nodes/WebGPUNodeBuilder.js b/examples/jsm/renderers/webgpu/nodes/WebGPUNodeBuilder.js
index 3a0951fd2f7493..ce9e41554c7c5b 100644
--- a/examples/jsm/renderers/webgpu/nodes/WebGPUNodeBuilder.js
+++ b/examples/jsm/renderers/webgpu/nodes/WebGPUNodeBuilder.js
@@ -1,4 +1,4 @@
-import WebGPUNodeUniformsGroup from './WebGPUNodeUniformsGroup.js';
+import WebGPUUniformsGroup from '../WebGPUUniformsGroup.js';
 import {
 	FloatNodeUniform, Vector2NodeUniform, Vector3NodeUniform, Vector4NodeUniform,
 	ColorNodeUniform, Matrix3NodeUniform, Matrix4NodeUniform
@@ -7,6 +7,7 @@ import WebGPUNodeSampler from './WebGPUNodeSampler.js';
 import { WebGPUNodeSampledTexture, WebGPUNodeSampledCubeTexture } from './WebGPUNodeSampledTexture.js';
 
 import WebGPUUniformBuffer from '../WebGPUUniformBuffer.js';
+import WebGPUStorageBuffer from '../WebGPUStorageBuffer.js';
 import { getVectorLength, getStrideLength } from '../WebGPUBufferUtils.js';
 
 import NodeBuilder from 'three-nodes/core/NodeBuilder.js';
@@ -16,6 +17,12 @@ import CodeNode from 'three-nodes/core/CodeNode.js';
 
 import { NodeMaterial } from 'three-nodes/materials/Materials.js';
 
+const gpuShaderStageLib = {
+	'vertex': GPUShaderStage.VERTEX,
+	'fragment': GPUShaderStage.FRAGMENT,
+	'compute': GPUShaderStage.COMPUTE
+};
+
 const supports = {
 	instance: true
 };
@@ -93,8 +100,8 @@ class WebGPUNodeBuilder extends NodeBuilder {
 		this.lightNode = null;
 		this.fogNode = null;
 
-		this.bindings = { vertex: [], fragment: [] };
-		this.bindingsOffset = { vertex: 0, fragment: 0 };
+		this.bindings = { vertex: [], fragment: [], compute: [] };
+		this.bindingsOffset = { vertex: 0, fragment: 0, compute: 0 };
 
 		this.uniformsGroup = {};
 
@@ -104,7 +111,17 @@ class WebGPUNodeBuilder extends NodeBuilder {
 
 	build() {
 
-		NodeMaterial.fromMaterial( this.material ).build( this );
+		const { object, material } = this;
+
+		if ( material !== null ) {
+
+			NodeMaterial.fromMaterial( material ).build( this );
+
+		} else {
+
+			this.addFlow( 'compute', object );
+
+		}
 
 		return super.build();
 
@@ -201,7 +218,7 @@ class WebGPUNodeBuilder extends NodeBuilder {
 
 				return name;
 
-			} else if ( type === 'buffer' ) {
+			} else if ( type === 'buffer' || type === 'storageBuffer' ) {
 
 				return `NodeBuffer_${node.node.id}.${name}`;
 
@@ -221,7 +238,7 @@ class WebGPUNodeBuilder extends NodeBuilder {
 
 		const bindings = this.bindings;
 
-		return [ ...bindings.vertex, ...bindings.fragment ];
+		return this.material !== null ? [ ...bindings.vertex, ...bindings.fragment ] : bindings.compute;
 
 	}
 
@@ -270,9 +287,11 @@ class WebGPUNodeBuilder extends NodeBuilder {
 
 				}
 
-			} else if ( type === 'buffer' ) {
+			} else if ( type === 'buffer' || type === 'storageBuffer' ) {
 
-				const buffer = new WebGPUUniformBuffer( 'NodeBuffer_' + node.id, node.value );
+				const bufferClass = type === 'storageBuffer' ? WebGPUStorageBuffer : WebGPUUniformBuffer;
+				const buffer = new bufferClass( 'NodeBuffer_' + node.id, node.value );
+				buffer.setVisibility( gpuShaderStageLib[ shaderStage ] );
 
 				// add first textures in sequence and group for last
 				const lastBinding = bindings[ bindings.length - 1 ];
@@ -288,7 +307,8 @@ class WebGPUNodeBuilder extends NodeBuilder {
 
 				if ( uniformsGroup === undefined ) {
 
-					uniformsGroup = new WebGPUNodeUniformsGroup( shaderStage );
+					uniformsGroup = new WebGPUUniformsGroup( 'nodeUniforms' );
+					uniformsGroup.setVisibility( gpuShaderStageLib[ shaderStage ] );
 
 					this.uniformsGroup[ shaderStage ] = uniformsGroup;
 
@@ -348,11 +368,7 @@ class WebGPUNodeBuilder extends NodeBuilder {
 
 		this.builtins.add( 'instance_index' );
 
-		if ( shaderStage === 'vertex' ) {
-
-			return 'instanceIndex';
-
-		}
+		return 'instanceIndex';
 
 	}
 
@@ -360,9 +376,13 @@ class WebGPUNodeBuilder extends NodeBuilder {
 
 		const snippets = [];
 
-		if ( shaderStage === 'vertex' ) {
+		if ( shaderStage === 'vertex' || shaderStage === 'compute' ) {
+
+			if ( shaderStage === 'compute' ) {
 
-			if ( this.builtins.has( 'instance_index' ) ) {
+				snippets.push( `@builtin( global_invocation_id ) id : vec3<u32>` );
+
+			} else if ( this.builtins.has( 'instance_index' ) ) {
 
 				snippets.push( `@builtin( instance_index ) instanceIndex : u32` );
 
@@ -477,15 +497,17 @@ class WebGPUNodeBuilder extends NodeBuilder {
 
 				bindingSnippets.push( `@group( 0 ) @binding( ${index ++} ) var ${uniform.name} : texture_cube<f32>;` );
 
-			} else if ( uniform.type === 'buffer' ) {
+			} else if ( uniform.type === 'buffer' || uniform.type === 'storageBuffer' ) {
 
 				const bufferNode = uniform.node;
 				const bufferType = this.getType( bufferNode.bufferType );
 				const bufferCount = bufferNode.bufferCount;
 
-				const bufferSnippet = `\t${uniform.name} : array< ${bufferType}, ${bufferCount} >\n`;
+				const bufferCountSnippet = bufferCount > 0 ? ', ' + bufferCount : '';
+				const bufferSnippet = `\t${uniform.name} : array< ${bufferType}${bufferCountSnippet} >\n`;
+				const bufferAccessMode = bufferNode.isStorageBufferNode ? 'storage,read_write' : 'uniform';
 
-				bufferSnippets.push( this._getWGSLUniforms( 'NodeBuffer_' + bufferNode.id, bufferSnippet, index ++ ) );
+				bufferSnippets.push( this._getWGSLStructBinding( 'NodeBuffer_' + bufferNode.id, bufferSnippet, bufferAccessMode, index ++ ) );
 
 			} else {
 
@@ -512,7 +534,7 @@ class WebGPUNodeBuilder extends NodeBuilder {
 
 		if ( groupSnippets.length > 0 ) {
 
-			code += this._getWGSLUniforms( 'NodeUniforms', groupSnippets.join( ',\n' ), index ++ );
+			code += this._getWGSLStructBinding( 'NodeUniforms', groupSnippets.join( ',\n' ), 'uniform', index ++ );
 
 		}
 
@@ -522,7 +544,7 @@ class WebGPUNodeBuilder extends NodeBuilder {
 
 	buildCode() {
 
-		const shadersData = { fragment: {}, vertex: {} };
+		const shadersData = this.material !== null ? { fragment: {}, vertex: {} } : { compute: {} };
 
 		for ( const shaderStage in shadersData ) {
 
@@ -548,7 +570,7 @@ class WebGPUNodeBuilder extends NodeBuilder {
 
 				flow += `${ flowSlotData.code }\n\t`;
 
-				if ( node === mainNode ) {
+				if ( node === mainNode && shaderStage !== 'compute' ) {
 
 					flow += '// FLOW RESULT\n\t';
 
@@ -579,8 +601,16 @@ class WebGPUNodeBuilder extends NodeBuilder {
 
 		}
 
-		this.vertexShader = this._getWGSLVertexCode( shadersData.vertex );
-		this.fragmentShader = this._getWGSLFragmentCode( shadersData.fragment );
+		if ( this.material !== null ) {
+
+			this.vertexShader = this._getWGSLVertexCode( shadersData.vertex );
+			this.fragmentShader = this._getWGSLFragmentCode( shadersData.fragment );
+
+		} else {
+
+			this.computeShader = this._getWGSLComputeCode( shadersData.compute, ( this.object.workgroupSize || [ 64 ] ).join( ', ' ) );
+
+		}
 
 	}
 
@@ -679,6 +709,35 @@ fn main( ${shaderData.varys} ) -> @location( 0 ) vec4<f32> {
 	// flow
 	${shaderData.flow}
 
+}
+`;
+
+	}
+
+	_getWGSLComputeCode( shaderData, workgroupSize ) {
+
+		return `${ this.getSignature() }
+// system
+var<private> instanceIndex : u32;
+
+// uniforms
+${shaderData.uniforms}
+
+// codes
+${shaderData.codes}
+
+@stage( compute ) @workgroup_size( ${workgroupSize} )
+fn main( ${shaderData.attributes} ) {
+
+	// system
+	instanceIndex = id.x * 3u;
+
+	// vars
+	${shaderData.vars}
+
+	// flow
+	${shaderData.flow}
+
 }
 `;
 
@@ -693,14 +752,14 @@ ${vars}
 
 	}
 
-	_getWGSLUniforms( name, vars, binding = 0, group = 0 ) {
+	_getWGSLStructBinding( name, vars, access, binding = 0, group = 0 ) {
 
 		const structName = name + 'Struct';
 		const structSnippet = this._getWGSLStruct( structName, vars );
 
 		return `${structSnippet}
 @binding( ${binding} ) @group( ${group} )
-var<uniform> ${name} : ${structName};`;
+var<${access}> ${name} : ${structName};`;
 
 	}
 
diff --git a/examples/jsm/renderers/webgpu/nodes/WebGPUNodeUniformsGroup.js b/examples/jsm/renderers/webgpu/nodes/WebGPUNodeUniformsGroup.js
deleted file mode 100644
index 793c725f2ce411..00000000000000
--- a/examples/jsm/renderers/webgpu/nodes/WebGPUNodeUniformsGroup.js
+++ /dev/null
@@ -1,20 +0,0 @@
-import WebGPUUniformsGroup from '../WebGPUUniformsGroup.js';
-
-class WebGPUNodeUniformsGroup extends WebGPUUniformsGroup {
-
-	constructor( shaderStage ) {
-
-		super( 'nodeUniforms' );
-
-		let shaderStageVisibility;
-
-		if ( shaderStage === 'vertex' ) shaderStageVisibility = GPUShaderStage.VERTEX;
-		else if ( shaderStage === 'fragment' ) shaderStageVisibility = GPUShaderStage.FRAGMENT;
-
-		this.setVisibility( shaderStageVisibility );
-
-	}
-
-}
-
-export default WebGPUNodeUniformsGroup;

From 7db9a5a217d89135ee6dd28f606cdfae966c281b Mon Sep 17 00:00:00 2001
From: sunag <jcdeconto@igalia.com>
Date: Sat, 16 Apr 2022 14:21:19 -0300
Subject: [PATCH 06/11] update webgpu_compute example

---
 examples/webgpu_compute.html | 151 ++++++++++++-----------------------
 1 file changed, 50 insertions(+), 101 deletions(-)

diff --git a/examples/webgpu_compute.html b/examples/webgpu_compute.html
index 9b92cefb2bd20d..94bc22a60e8cb6 100644
--- a/examples/webgpu_compute.html
+++ b/examples/webgpu_compute.html
@@ -28,23 +28,22 @@
 			import * as THREE from 'three';
 			import * as Nodes from 'three-nodes/Nodes.js';
 
+			import {
+				compute,
+				color, add, uniform, element, storage, func,
+				positionLocal, instanceIndex
+			} from 'three-nodes/Nodes.js';
+
 			import { GUI } from './jsm/libs/lil-gui.module.min.js';
 
 			import WebGPU from './jsm/capabilities/WebGPU.js';
 			import WebGPURenderer from './jsm/renderers/webgpu/WebGPURenderer.js';
 
-			import WebGPUStorageBuffer from './jsm/renderers/webgpu/WebGPUStorageBuffer.js';
-			import WebGPUUniformBuffer from './jsm/renderers/webgpu/WebGPUUniformBuffer.js';
-			import * as WebGPUBufferUtils from './jsm/renderers/webgpu/WebGPUBufferUtils.js';
-			import WebGPUUniformsGroup from './jsm/renderers/webgpu/WebGPUUniformsGroup.js';
-			import { Vector2Uniform } from './jsm/renderers/webgpu/WebGPUUniform.js';
-
 			let camera, scene, renderer;
-			let pointer;
-			let scaleUniformBuffer;
-			const scaleVector = new THREE.Vector3( 1, 1, 1 );
+			let computeNode;
 
-			const computeParams = [];
+			const pointer = new THREE.Vector2( - 10.0, - 10.0 ); // Out of bounds first
+			const scaleVector = new THREE.Vector2( 1, 1 );
 
 			init().then( animate ).catch( error );
 
@@ -62,10 +61,11 @@
 				camera.position.z = 1;
 
 				scene = new THREE.Scene();
-				scene.background = new THREE.Color( 0x000000 );
+
+				// initialize particles
 
 				const particleNum = 65000; // 16-bit limit
-				const particleSize = 4; // 16-byte stride align
+				const particleSize = 3; // vec3
 
 				const particleArray = new Float32Array( particleNum * particleSize );
 				const velocityArray = new Float32Array( particleNum * particleSize );
@@ -74,81 +74,25 @@
 
 					const r = Math.random() * 0.01 + 0.0005;
 					const degree = Math.random() * 360;
+
 					velocityArray[ i + 0 ] = r * Math.sin( degree * Math.PI / 180 );
 					velocityArray[ i + 1 ] = r * Math.cos( degree * Math.PI / 180 );
 
 				}
 
-				const particleBuffer = new WebGPUStorageBuffer( 'particle', new THREE.BufferAttribute( particleArray, particleSize ) );
-				const velocityBuffer = new WebGPUStorageBuffer( 'velocity', new THREE.BufferAttribute( velocityArray, particleSize ) );
-
-				const scaleUniformLength = WebGPUBufferUtils.getVectorLength( 2, 3 ); // two vector3 for array
-
-				scaleUniformBuffer = new WebGPUUniformBuffer( 'scaleUniform', new Float32Array( scaleUniformLength ) );
-
-				pointer = new THREE.Vector2( - 10.0, - 10.0 ); // Out of bounds first
-
-				const pointerGroup = new WebGPUUniformsGroup( 'mouseUniforms' ).addUniform(
-					new Vector2Uniform( 'pointer', pointer )
-				);
-
-				// Object keys need follow the binding shader sequence
-
-				const computeBindings = [
-					particleBuffer,
-					velocityBuffer,
-					scaleUniformBuffer,
-					pointerGroup
-				];
-
-				const computeShader = `
-
-					//
-					// Buffer
-					//
-
-					struct Particle {
-						value : array< vec4<f32> >
-					};
-					@binding( 0 ) @group( 0 )
-					var<storage,read_write> particle : Particle;
-
-					struct Velocity {
-						value : array< vec4<f32> >
-					};
-					@binding( 1 ) @group( 0 )
-					var<storage,read_write> velocity : Velocity;
+				// create buffers
 
-					//
-					// Uniforms
-					//
+				const particleBuffer = new THREE.BufferAttribute( particleArray, particleSize );
+				const velocityBuffer = new THREE.BufferAttribute( velocityArray, particleSize );
 
-					struct Scale {
-						value : array< vec3<f32>, 2 >
-					};
-					@binding( 2 ) @group( 0 )
-					var<uniform> scaleUniform : Scale;
+				const particleBufferNode = storage( particleBuffer, 'vec3' );
+				const velocityBufferNode = storage( velocityBuffer, 'vec3' );
 
-					struct MouseUniforms {
-						pointer : vec2<f32>
-					};
-					@binding( 3 ) @group( 0 )
-					var<uniform> mouseUniforms : MouseUniforms;
+				// create wgsl function
 
-					@stage( compute ) @workgroup_size( 64 )
-					fn main( @builtin(global_invocation_id) id : vec3<u32> ) {
+				const WGSLFnNode = func( `( pointer:vec2<f32>, limit:vec2<f32> ) -> vec3<f32> {
 
-						// get particle index
-
-						let index : u32 = id.x * 3u;
-
-						// update speed
-
-						var position : vec4<f32> = particle.value[ index ] + velocity.value[ index ];
-
-						// update limit
-
-						let limit : vec2<f32> = scaleUniform.value[ 0 ].xy;
+						var position = particle + velocity;
 
 						if ( abs( position.x ) >= limit.x ) {
 
@@ -162,7 +106,7 @@
 
 							}
 
-							velocity.value[ index ].x = - velocity.value[ index ].x;
+							velocity.x = - velocity.x;
 
 						}
 
@@ -178,17 +122,15 @@
 
 							}
 
-							velocity.value[ index ].y = - velocity.value[ index ].y;
+							velocity.y = - velocity.y ;
 
 						}
 
-						// update mouse
-
-						let POINTER_SIZE : f32 = .1;
+						let POINTER_SIZE = .1;
 
-						let dx : f32 = mouseUniforms.pointer.x - position.x;
-						let dy : f32 = mouseUniforms.pointer.y - position.y;
-						let distanceFromPointer : f32 = sqrt( dx * dx + dy * dy );
+						let dx = pointer.x - position.x;
+						let dy = pointer.y - position.y;
+						let distanceFromPointer = sqrt( dx * dx + dy * dy );
 
 						if ( distanceFromPointer <= POINTER_SIZE ) {
 
@@ -198,28 +140,37 @@
 
 						}
 
-						// update buffer
-
-						particle.value[ index ] = position;
+						return position;
 
 					}
+				` );
 
-				`;
+				// define particle and velocity keywords in wgsl function
 
-				computeParams.push( {
-					num: particleNum,
-					shader: computeShader,
-					bindings: computeBindings
+				WGSLFnNode.keywords[ 'particle' ] = element( particleBufferNode, instanceIndex );
+				WGSLFnNode.keywords[ 'velocity' ] = element( velocityBufferNode, instanceIndex );
+
+				// create a call function node
+
+				const WGSLFnCallNode = WGSLFnNode.call( {
+					pointer: uniform( pointer ),
+					limit: uniform( scaleVector )
 				} );
 
-				// Use a compute shader to animate the point cloud's vertex data.
+				// compute
 
-				const pointsGeometry = new THREE.BufferGeometry().setAttribute(
-					'position', particleBuffer.attribute
-				);
+				computeNode = compute( particleNum );
+
+				// WGSLFnCallNode can be other kind of node, like a ShaderNode for example
+				computeNode.assign( particleBufferNode, WGSLFnCallNode );
+
+				// use a compute shader to animate the point cloud's vertex data.
+
+				const pointsGeometry = new THREE.BufferGeometry();
+				pointsGeometry.setAttribute( 'position', particleBuffer );
 
 				const pointsMaterial = new Nodes.PointsNodeMaterial();
-				pointsMaterial.colorNode = new Nodes.OperatorNode( '+', new Nodes.PositionNode(), new Nodes.UniformNode( new THREE.Color( 0xFFFFFF ) ) );
+				pointsMaterial.colorNode = add( positionLocal, color( 0xFFFFFF ) );
 
 				const mesh = new THREE.Points( pointsGeometry, pointsMaterial );
 				scene.add( mesh );
@@ -270,11 +221,9 @@
 
 				requestAnimationFrame( animate );
 
-				renderer.compute( computeParams );
+				renderer.compute( computeNode );
 				renderer.render( scene, camera );
 
-				scaleVector.toArray( scaleUniformBuffer.buffer, 0 );
-
 			}
 
 			function error( error ) {

From c8570accb5fcd4ca58b699c205b56d6a540a03bc Mon Sep 17 00:00:00 2001
From: sunag <jcdeconto@igalia.com>
Date: Sat, 16 Apr 2022 14:42:43 -0300
Subject: [PATCH 07/11] cleanup

---
 examples/webgpu_compute.html | 5 +++--
 1 file changed, 3 insertions(+), 2 deletions(-)

diff --git a/examples/webgpu_compute.html b/examples/webgpu_compute.html
index 94bc22a60e8cb6..ea5f167196e85e 100644
--- a/examples/webgpu_compute.html
+++ b/examples/webgpu_compute.html
@@ -90,7 +90,7 @@
 
 				// create wgsl function
 
-				const WGSLFnNode = func( `( pointer:vec2<f32>, limit:vec2<f32> ) -> vec3<f32> {
+				const WGSLFnNode = func( `( particle:vec3<f32>, pointer:vec2<f32>, limit:vec2<f32> ) -> vec3<f32> {
 
 						var position = particle + velocity;
 
@@ -146,13 +146,14 @@
 				` );
 
 				// define particle and velocity keywords in wgsl function
+				// it's used in case of needed change a global variable like this storageBuffer
 
-				WGSLFnNode.keywords[ 'particle' ] = element( particleBufferNode, instanceIndex );
 				WGSLFnNode.keywords[ 'velocity' ] = element( velocityBufferNode, instanceIndex );
 
 				// create a call function node
 
 				const WGSLFnCallNode = WGSLFnNode.call( {
+					particle: element( particleBufferNode, instanceIndex ),
 					pointer: uniform( pointer ),
 					limit: uniform( scaleVector )
 				} );

From 69958226fec7f4fe233caa4277f09ca8eba9c62a Mon Sep 17 00:00:00 2001
From: sunag <jcdeconto@igalia.com>
Date: Thu, 21 Apr 2022 16:51:54 -0300
Subject: [PATCH 08/11] fix ToneMappingNode in non-physical material

---
 examples/jsm/nodes/materials/MeshStandardNodeMaterial.js | 4 +++-
 1 file changed, 3 insertions(+), 1 deletion(-)

diff --git a/examples/jsm/nodes/materials/MeshStandardNodeMaterial.js b/examples/jsm/nodes/materials/MeshStandardNodeMaterial.js
index 9a1b6ec1e3252f..332f8f62d6a8cd 100644
--- a/examples/jsm/nodes/materials/MeshStandardNodeMaterial.js
+++ b/examples/jsm/nodes/materials/MeshStandardNodeMaterial.js
@@ -63,7 +63,9 @@ export default class MeshStandardNodeMaterial extends NodeMaterial {
 
 		const outgoingLightNode = super.generateLight( builder, { diffuseColorNode, lightNode, lightingModelNode: PhysicalLightingModel } );
 
-		// @TODO: add IBL code here
+		// TONE MAPPING
+
+		if ( renderer.toneMappingNode ) outgoingLightNode = context( renderer.toneMappingNode, { color: outgoingLightNode } );
 
 		return outgoingLightNode;
 

From 59bb988409c8de9aa0986dfca3a4931016ebaca6 Mon Sep 17 00:00:00 2001
From: sunag <jcdeconto@igalia.com>
Date: Thu, 21 Apr 2022 19:04:56 -0300
Subject: [PATCH 09/11] fixes and cleanup

---
 examples/jsm/nodes/core/NodeBuilder.js        | 67 +++++++------------
 examples/jsm/nodes/gpgpu/ComputeNode.js       | 27 ++------
 .../materials/MeshStandardNodeMaterial.js     |  6 +-
 examples/jsm/nodes/materials/NodeMaterial.js  |  8 +--
 .../jsm/nodes/shadernode/ShaderNodeUtils.js   | 23 +++++--
 examples/webgpu_compute.html                  | 34 +++++++---
 examples/webgpu_nodes_playground.html         |  1 -
 7 files changed, 78 insertions(+), 88 deletions(-)

diff --git a/examples/jsm/nodes/core/NodeBuilder.js b/examples/jsm/nodes/core/NodeBuilder.js
index 02fe433b42cfdf..6b05a622b5cbdb 100644
--- a/examples/jsm/nodes/core/NodeBuilder.js
+++ b/examples/jsm/nodes/core/NodeBuilder.js
@@ -8,9 +8,17 @@ import { NodeUpdateType } from './constants.js';
 
 import { REVISION, LinearEncoding } from 'three';
 
-export const shaderStages = [ 'fragment', 'vertex' ];
+export const shaderStages = [ 'fragment', 'vertex', 'compute' ];
 export const vector = [ 'x', 'y', 'z', 'w' ];
 
+const typeFromLength = new Map();
+typeFromLength.set( 1, 'float' );
+typeFromLength.set( 2, 'vec2' );
+typeFromLength.set( 3, 'vec3' );
+typeFromLength.set( 4, 'vec4' );
+typeFromLength.set( 9, 'mat3' );
+typeFromLength.set( 16, 'mat4' );
+
 const toFloat = ( value ) => {
 
 	value = Number( value );
@@ -331,16 +339,9 @@ class NodeBuilder {
 
 	}
 
-	getTypeFromLength( type ) {
-
-		if ( type === 1 ) return 'float';
-		if ( type === 2 ) return 'vec2';
-		if ( type === 3 ) return 'vec3';
-		if ( type === 4 ) return 'vec4';
-		if ( type === 9 ) return 'mat3';
-		if ( type === 16 ) return 'mat4';
+	getTypeFromLength( length ) {
 
-		return 0;
+		return typeFromLength.get( length );
 
 	}
 
@@ -617,51 +618,33 @@ class NodeBuilder {
 
 	build() {
 
-		if ( this.material !== null ) {
-
-			// stage 1: analyze nodes to possible optimization and validation
-
-			for ( const shaderStage of shaderStages ) {
-
-				this.setShaderStage( shaderStage );
+		// stage 1: analyze nodes to possible optimization and validation
 
-				const flowNodes = this.flowNodes[ shaderStage ];
+		for ( const shaderStage of shaderStages ) {
 
-				for ( const node of flowNodes ) {
-
-					node.analyze( this );
-
-				}
-
-			}
+			this.setShaderStage( shaderStage );
 
-			// stage 2: pre-build vertex code used in fragment shader
+			const flowNodes = this.flowNodes[ shaderStage ];
 
-			if ( this.context.vertex && this.context.vertex.isNode ) {
+			for ( const node of flowNodes ) {
 
-				this.flowNodeFromShaderStage( 'vertex', this.context.vertex );
+				node.analyze( this );
 
 			}
 
-			// stage 3: generate shader
-
-			for ( const shaderStage of shaderStages ) {
-
-				this.setShaderStage( shaderStage );
-
-				const flowNodes = this.flowNodes[ shaderStage ];
+		}
 
-				for ( const node of flowNodes ) {
+		// stage 2: pre-build vertex code used in fragment shader
 
-					this.flowNode( node, shaderStage );
+		if ( this.context.vertex && this.context.vertex.isNode ) {
 
-				}
+			this.flowNodeFromShaderStage( 'vertex', this.context.vertex );
 
-			}
+		}
 
-		} else {
+		// stage 3: generate shader
 
-			const shaderStage = 'compute';
+		for ( const shaderStage of shaderStages ) {
 
 			this.setShaderStage( shaderStage );
 
@@ -669,8 +652,6 @@ class NodeBuilder {
 
 			for ( const node of flowNodes ) {
 
-				node.analyze( this );
-
 				this.flowNode( node, shaderStage );
 
 			}
diff --git a/examples/jsm/nodes/gpgpu/ComputeNode.js b/examples/jsm/nodes/gpgpu/ComputeNode.js
index 4718613e685807..4da4469476353f 100644
--- a/examples/jsm/nodes/gpgpu/ComputeNode.js
+++ b/examples/jsm/nodes/gpgpu/ComputeNode.js
@@ -1,5 +1,4 @@
 import Node from '../core/Node.js';
-import { assign, element, instanceIndex } from '../shadernode/ShaderNodeElements.js';
 import { NodeUpdateType } from '../core/constants.js';
 
 class ComputeNode extends Node {
@@ -13,23 +12,7 @@ class ComputeNode extends Node {
 		this.dispatchCount = dispatchCount;
 		this.workgroupSize = workgroupSize;
 
-		this.assigns = [];
-
-	}
-
-	getMainStorageBufferNode() {
-
-		const assigns = this.assigns;
-
-		return assigns[ assigns.length - 1 ].storageBufferNode;
-
-	}
-
-	assign( storageBufferNode, sourceNode ) {
-
-		this.assigns.push( { storageBufferNode, sourceNode } );
-
-		return this;
+		this.computeNode = null;
 
 	}
 
@@ -41,13 +24,15 @@ class ComputeNode extends Node {
 
 	generate( builder ) {
 
-		const { renderer, shaderStage } = builder;
+		const { shaderStage } = builder;
 
 		if ( shaderStage === 'compute' ) {
 
-			for ( const { storageBufferNode, sourceNode } of this.assigns ) {
+			const snippet = this.computeNode.build( builder, 'void' );
+
+			if ( snippet !== '' ) {
 
-				assign( element( storageBufferNode, instanceIndex ), sourceNode ).build( builder );
+				builder.addFlowCode( snippet );
 
 			}
 
diff --git a/examples/jsm/nodes/materials/MeshStandardNodeMaterial.js b/examples/jsm/nodes/materials/MeshStandardNodeMaterial.js
index 332f8f62d6a8cd..34374ce1e68dff 100644
--- a/examples/jsm/nodes/materials/MeshStandardNodeMaterial.js
+++ b/examples/jsm/nodes/materials/MeshStandardNodeMaterial.js
@@ -1,7 +1,7 @@
 import NodeMaterial from './NodeMaterial.js';
 import {
 	float, vec3, vec4,
-	assign, label, mul, invert, mix,
+	context, assign, label, mul, invert, mix,
 	normalView,
 	materialRoughness, materialMetalness
 } from '../shadernode/ShaderNodeElements.js';
@@ -61,10 +61,12 @@ export default class MeshStandardNodeMaterial extends NodeMaterial {
 
 	generateLight( builder, { diffuseColorNode, lightNode } ) {
 
-		const outgoingLightNode = super.generateLight( builder, { diffuseColorNode, lightNode, lightingModelNode: PhysicalLightingModel } );
+		let outgoingLightNode = super.generateLight( builder, { diffuseColorNode, lightNode, lightingModelNode: PhysicalLightingModel } );
 
 		// TONE MAPPING
 
+		const renderer = builder.renderer;
+
 		if ( renderer.toneMappingNode ) outgoingLightNode = context( renderer.toneMappingNode, { color: outgoingLightNode } );
 
 		return outgoingLightNode;
diff --git a/examples/jsm/nodes/materials/NodeMaterial.js b/examples/jsm/nodes/materials/NodeMaterial.js
index eaf220c2a61e80..7f51110c62b782 100644
--- a/examples/jsm/nodes/materials/NodeMaterial.js
+++ b/examples/jsm/nodes/materials/NodeMaterial.js
@@ -4,7 +4,7 @@ import ExpressionNode from '../core/ExpressionNode.js';
 import {
 	float, vec3, vec4,
 	assign, label, mul, add, bypass,
-	positionLocal, skinning, instance, modelViewProjection, context, lightContext, colorSpace,
+	positionLocal, skinning, instance, modelViewProjection, lightContext, colorSpace,
 	materialAlphaTest, materialColor, materialOpacity
 } from '../shadernode/ShaderNodeElements.js';
 
@@ -112,16 +112,10 @@ class NodeMaterial extends ShaderMaterial {
 
 	generateOutput( builder, { diffuseColorNode, outgoingLightNode } ) {
 
-		const { renderer } = builder;
-
 		// OUTPUT
 
 		let outputNode = vec4( outgoingLightNode, diffuseColorNode.a );
 
-		// TONE MAPPING
-
-		if ( renderer.toneMappingNode ) outputNode = context( renderer.toneMappingNode, { color: outputNode } );
-
 		// ENCODING
 
 		outputNode = colorSpace( outputNode, builder.renderer.outputEncoding );
diff --git a/examples/jsm/nodes/shadernode/ShaderNodeUtils.js b/examples/jsm/nodes/shadernode/ShaderNodeUtils.js
index 81d492e024f0a8..45a5a0c2cc183e 100644
--- a/examples/jsm/nodes/shadernode/ShaderNodeUtils.js
+++ b/examples/jsm/nodes/shadernode/ShaderNodeUtils.js
@@ -141,13 +141,28 @@ const ShaderNodeProxy = function ( NodeClass, scope = null, factor = null ) {
 
 export const ShaderNodeScript = function ( jsFunc ) {
 
-	return { call: ( inputs, builder ) => {
+	//@TODO: Move this to Node extended class
 
-		inputs = new ShaderNodeObjects( inputs );
+	const self =
+	{
+		build: ( builder ) => {
 
-		return new ShaderNodeObject( jsFunc( inputs, builder ) );
+			self.call( {}, builder );
 
-	} };
+			return '';
+
+		},
+
+		call: ( inputs, builder ) => {
+
+			inputs = new ShaderNodeObjects( inputs );
+
+			return new ShaderNodeObject( jsFunc( inputs, builder ) );
+
+		}
+	};
+
+	return self;
 
 };
 
diff --git a/examples/webgpu_compute.html b/examples/webgpu_compute.html
index ea5f167196e85e..12c96d40ce0464 100644
--- a/examples/webgpu_compute.html
+++ b/examples/webgpu_compute.html
@@ -31,6 +31,7 @@
 			import {
 				compute,
 				color, add, uniform, element, storage, func,
+				assign, float, mul,
 				positionLocal, instanceIndex
 			} from 'three-nodes/Nodes.js';
 
@@ -90,7 +91,7 @@
 
 				// create wgsl function
 
-				const WGSLFnNode = func( `( particle:vec3<f32>, pointer:vec2<f32>, limit:vec2<f32> ) -> vec3<f32> {
+				const WGSLFnNode = func( `( pointer:vec2<f32>, limit:vec2<f32> ) {
 
 						var position = particle + velocity;
 
@@ -140,7 +141,7 @@
 
 						}
 
-						return position;
+						particle = position;
 
 					}
 				` );
@@ -148,22 +149,35 @@
 				// define particle and velocity keywords in wgsl function
 				// it's used in case of needed change a global variable like this storageBuffer
 
-				WGSLFnNode.keywords[ 'velocity' ] = element( velocityBufferNode, instanceIndex );
+				const particleNode = element( particleBufferNode, instanceIndex );
+				const velocityNode = element( velocityBufferNode, instanceIndex );
 
-				// create a call function node
+				WGSLFnNode.keywords[ 'particle' ] = particleNode;
+				WGSLFnNode.keywords[ 'velocity' ] = velocityNode;
 
-				const WGSLFnCallNode = WGSLFnNode.call( {
-					particle: element( particleBufferNode, instanceIndex ),
+				// compute
+
+				computeNode = compute( particleNum );
+
+				// Example 1: Calling a WGSL function
+
+				computeNode.computeNode = WGSLFnNode.call( {
 					pointer: uniform( pointer ),
 					limit: uniform( scaleVector )
 				} );
 
-				// compute
+				// Example 2: Creating single storage assign
 
-				computeNode = compute( particleNum );
+				//computeNode.computeNode = assign( particleNode, add( particleNode, velocityNode ) );
+
+				// Example 3: Creating multiples storage assign
+
+				/*computeNode.computeNode = new Nodes.ShaderNode( ( {}, builder ) => {
+
+					assign( particleNode, add( particleNode, velocityNode ) ).build( builder );
+					assign( velocityNode, mul( velocityNode, float( 0.99 ) ) ).build( builder );
 
-				// WGSLFnCallNode can be other kind of node, like a ShaderNode for example
-				computeNode.assign( particleBufferNode, WGSLFnCallNode );
+				} );/**/
 
 				// use a compute shader to animate the point cloud's vertex data.
 
diff --git a/examples/webgpu_nodes_playground.html b/examples/webgpu_nodes_playground.html
index 228be5a0921ec0..3d7792b0c2e3c5 100644
--- a/examples/webgpu_nodes_playground.html
+++ b/examples/webgpu_nodes_playground.html
@@ -73,7 +73,6 @@
 			let stats;
 			let camera, scene, renderer;
 			let model;
-			let nodeLights;
 
 			init().then( animate ).catch( error => console.error( error ) );
 

From 41898a69174146e258cfcd851600fc4b661254a7 Mon Sep 17 00:00:00 2001
From: sunag <jcdeconto@igalia.com>
Date: Thu, 21 Apr 2022 19:37:52 -0300
Subject: [PATCH 10/11] fixes and ignore tonemapping in non physical material

---
 examples/jsm/nodes/core/NodeBuilder.js             |  3 ++-
 .../jsm/renderers/webgl/nodes/WebGLNodeBuilder.js  | 14 ++++++++++----
 2 files changed, 12 insertions(+), 5 deletions(-)

diff --git a/examples/jsm/nodes/core/NodeBuilder.js b/examples/jsm/nodes/core/NodeBuilder.js
index 6b05a622b5cbdb..1a541c954b5ed8 100644
--- a/examples/jsm/nodes/core/NodeBuilder.js
+++ b/examples/jsm/nodes/core/NodeBuilder.js
@@ -8,7 +8,8 @@ import { NodeUpdateType } from './constants.js';
 
 import { REVISION, LinearEncoding } from 'three';
 
-export const shaderStages = [ 'fragment', 'vertex', 'compute' ];
+export const defaultShaderStages = [ 'fragment', 'vertex' ];
+export const shaderStages = [ ...defaultShaderStages, 'compute' ];
 export const vector = [ 'x', 'y', 'z', 'w' ];
 
 const typeFromLength = new Map();
diff --git a/examples/jsm/renderers/webgl/nodes/WebGLNodeBuilder.js b/examples/jsm/renderers/webgl/nodes/WebGLNodeBuilder.js
index ba4a75793367cb..7c1a3a8328ca15 100644
--- a/examples/jsm/renderers/webgl/nodes/WebGLNodeBuilder.js
+++ b/examples/jsm/renderers/webgl/nodes/WebGLNodeBuilder.js
@@ -1,4 +1,4 @@
-import NodeBuilder, { shaderStages } from 'three-nodes/core/NodeBuilder.js';
+import NodeBuilder, { defaultShaderStages } from 'three-nodes/core/NodeBuilder.js';
 import NodeFrame from 'three-nodes/core/NodeFrame.js';
 import SlotNode from './SlotNode.js';
 import GLSLNodeParser from 'three-nodes/parsers/GLSLNodeParser.js';
@@ -85,6 +85,12 @@ class WebGLNodeBuilder extends NodeBuilder {
 
 		}
 
+		if ( material.isMeshStandardNodeMaterial !== true ) {
+
+			this.replaceCode( 'fragment', getIncludeSnippet( 'tonemapping_fragment' ), '' );
+
+		}
+
 		// parse inputs
 
 		if ( material.colorNode && material.colorNode.isNode ) {
@@ -329,7 +335,7 @@ class WebGLNodeBuilder extends NodeBuilder {
 
 		const shaderData = {};
 
-		for ( const shaderStage of shaderStages ) {
+		for ( const shaderStage of defaultShaderStages ) {
 
 			const uniforms = this.getUniforms( shaderStage );
 			const attributes = this.getAttributes( shaderStage );
@@ -515,7 +521,7 @@ ${this.shader[ getShaderStageProperty( shaderStage ) ]}
 
 		}
 
-		for ( const shaderStage of shaderStages ) {
+		for ( const shaderStage of defaultShaderStages ) {
 
 			this.addCodeAfterSnippet(
 				shaderStage,
@@ -529,7 +535,7 @@ ${this.shader[ getShaderStageProperty( shaderStage ) ]}
 
 	_addUniforms() {
 
-		for ( const shaderStage of shaderStages ) {
+		for ( const shaderStage of defaultShaderStages ) {
 
 			// uniforms
 

From aceb1f9adbffafabdf0d2a34a6ddb09da2c4b1b9 Mon Sep 17 00:00:00 2001
From: sunag <jcdeconto@igalia.com>
Date: Thu, 21 Apr 2022 19:55:24 -0300
Subject: [PATCH 11/11] update snapshots

---
 examples/screenshots/webgpu_compute.jpg       | Bin 6702 -> 3413 bytes
 examples/screenshots/webgpu_depth_texture.jpg | Bin 6702 -> 3413 bytes
 examples/screenshots/webgpu_instance_mesh.jpg | Bin 6702 -> 3413 bytes
 .../screenshots/webgpu_instance_uniform.jpg   | Bin 6784 -> 3413 bytes
 examples/screenshots/webgpu_lights_custom.jpg | Bin 6702 -> 3413 bytes
 .../screenshots/webgpu_lights_selective.jpg   | Bin 6702 -> 3413 bytes
 examples/screenshots/webgpu_materials.jpg     | Bin 6702 -> 3413 bytes
 .../screenshots/webgpu_nodes_playground.jpg   | Bin 6702 -> 3413 bytes
 examples/screenshots/webgpu_rtt.jpg           | Bin 6702 -> 3413 bytes
 examples/screenshots/webgpu_sandbox.jpg       | Bin 6702 -> 3413 bytes
 examples/screenshots/webgpu_skinning.jpg      | Bin 6702 -> 3413 bytes
 .../webgpu_skinning_instancing.jpg            | Bin 6702 -> 3413 bytes
 .../screenshots/webgpu_skinning_points.jpg    | Bin 6702 -> 3413 bytes
 13 files changed, 0 insertions(+), 0 deletions(-)

diff --git a/examples/screenshots/webgpu_compute.jpg b/examples/screenshots/webgpu_compute.jpg
index 35b06935974bb9787ab26daee5a8712bd8be7e9b..001fe203e0b94e476d12e5c1c76aa979a0fbc77f 100644
GIT binary patch
delta 21
dcmZ2ya#d=>CFV^++8mn`ghCiu{xkf)2>@O+2%`W1

literal 6702
zcmeHLX;c$iwyqE$M&JU<prXxHL2zJ@84+<rhK^i7KnQ_W!5AaZT$x7)RLljDDF}!t
zqYNP;B#{sZAR>dvWtKn~0>~T^1QLd%9_U;B=e-~Hd#n3TyK9|&s_L9t^?m2;v+Fwr
zeg=mD87p&3a{vMX0LaDxfc-!eAOhL6O)s)(qR>q*CMpUQ6&Dj1|5jk|E#fdZOk8}6
z#1{D0P1`t=l-Me{dGY-*5iuxKY%5G0_O1D!a$pT0y9HQ<Jb^-F0TEdUR2Bl(L$Wq7
zihYB&0ra~;M4+N#;;;<}5;7aGq0kMmVjD;|dWLVb0iv>E+x8wkCBEJEGHjo(oX(y2
zms|FqF0GZfYv-uwUitkle5-=uj-9(y52&dhJfx>@U}$7~?9AD7=I1Ret?Vy4I67T&
zcDd@|c@2FXbK{nue?VYRaLB!=`wt#QKf)y>CMDxjo~EW{W#{DP<-aN@d{b6VB2(U0
zQ0v~+zi;@^*woz7+11_C`?-%X!WtbL|1vQ-HO*c4y12B=<FBkj0O)r*8~%Ur*&q%P
z*&K7&CLf4M;6{PUii+($D!%QME$p)I_I)~cw#c21e_2`!->+-Ok-zeL`&I=Ny%AOJ
zCZ%tTevHuF{}ZEs5c&rnkO4?SAsZ70l?BWIVdNz6Lge;VK>cU+ZyEx#0xiBCZ`iPA
zDB6r;;iXg`Vah)`VtVU0pTA(^NkZaCKQ<M2q=01Ic>lGDB=@9pZX~A(1m3g>l$_py
z!2151k^F8z`RBd=xG@M*1_79kuq_f;mj(gRbNoJ5rdXIXCnBzZ0RnY~%U<WLIo$*r
zht1W*$C52ZB`X{~jgCLq|MG?OSo4o9av_mE{OA128Sh6R;0bjDf%c}8%K2TuZ;+E>
zfNzN5oGb^4tUwyAM`F;r(}(ll{Zh6YFRd{X*^(uZ9Mfbp>UTH=(NbYP-bliwghz@l
z521<ZpelRIx6Ouho5}!6|17z($TE~*h>0)>lL=m2^WZm)*dHJS@5xKv!IdJ<v^?62
z3t~n42u91DWGv)=w?ZgE9@?QY4YiPlJ}R?ASs^Fd%>SCZqln@u#E^J7Ls>T~SXww`
z>EnA@8Ctb&hq<v2QVX#-5cu3fjbd#NYiuO!KEiia5kD(nB!Gan(FeU08UrOSbeWkf
zp6@4=+(n!^u<Si#dZY@qINy2A<II{>#HW;9`t;~wzF;`NlU@I%LvY;T+U2LG?^@!_
zQc1LOPju~y!yO~>j11y<B%Zo@eQCkd(x~bEWKX89`NXJJorEdoxM3g{`Of!-&OlXp
zz&0#-#^huCdEEw?YSxTgsfE){_`ahfcXV<UH7_ppa;k(gCd*~B^;MHhudJ0y=P)*5
zaJoMpW42oepJ6;s&Yq0i5vCCuNK_%dWH5prFf8@uqgjr39Gaf)9l@b7Gq{{*WFZ9v
z)OaUqI|K$RQ~9jwMQav4#(nF$6f4}UHbRRlOwONX<k+=%=DQ9aYX0-RWM`AE=B?d_
z_t5H_KcBSso~wsy+FBv1_LZ!s1lYx76jK7+eBONB32T_Yr|6De(|SK7(X`xZfHLoF
z>79$qME?Z@bjGzoK+k%-dJ5KTU|><3Ctx-rEDuo0Dd|X}L$?Bb<odb}2t>(Ot8c{U
z`#<R;J1@l$SXy&I5j8R=>KaA^Q*PIk1YX_02U9+(@|@lBWqEBi@SvIoy~u~RlS-e1
zaMVt9W;*tq&<%1NWSGj`m^+dtc&pquk^6aNC2H)-sfA}TyR4n{D2_JG%cM28>A_b?
zof!9U+FWRVM;kjlwrL$ZJmU`Eczd`%i>YI}rH}BrCx;&65}uKy(>k%etmV)tsb?~$
zTDS4`<h~f4_bPa3HYdhW=(j?JOpFFb?Jvdpy+!tNHC{#6u|yJJI#;n0MOT!J_OXsc
z7P}_!(2a}X)FluY^W>fQs8~f=!n(7mYrVeGb5b19^T7xCqr)7kf3cIxy~sKQ_K!=d
z>ih6EgJIEMF4dLR_#D-~FoGXdNynmn(yDq(;%wADTjpyzR1tkNg=vJaM*lX3{c<3u
zi-H&a%AX%zGuE3RqfQ@=+r}tjq-Sb1ewC}fx($Bi5E6JK_jrq5m&&R?r44237ZNO&
z80jxW3H0?&99X~1599<|zwW8tH#2`3TQIGB<xT1l^t0n%c3#^OmK0uL;lV&c(r<{M
z+L60Rm(Py|2+2#~3UiZ*SEEoi4_86pHpM5aGUUKV4?&?B+Ps6HUr?<xsf}cksQAk2
zONeY6Z=OtUN7I8cG*0bMS#;B3iqdORm4#K(fLb<!nATo2%b1P6XMZ=z-^O98N^pF=
z*U!vdkW($s_w>%~&f&`7+H4iHe0_#vU4nAcJ=lX6<6Yi-!jU@_6LcEW&A*^pp2pD<
z=Z9>bS7{W6>KER8+h8BTuJroS<@9Opun$l7Y7Gb|(6)0GZL2$I2Yp{PX=YT;na`i1
zk>bWH-jYmIo~u}zg<_3XB>r9q0>14}oqXwetELXiAs~Qnv(qf0EeTUQe$)P@hWJq7
zFONYWNiOUb2vnd!z%Km2ZL=qben=)kSF1tb-WmG16je=-Bozt*=T(7-%x@V~4)Q_y
z_Rud<Elj;*oyXG#R*PTKoL(gde^K|0O&a)He`%~)yT!r@xBW9`x8=RaUiR@hFg6}(
z__4pT+w?A$VQd^;<UW?<IB2PLB3J7d;m09(m|C$jHcg5ieL2&nD^1v~)UaCqq_|Zy
zfPd7qJ$6*4KDt!r>7UZc<`-O*@+sFyz2~p-zT(2!CEKkqYVCP6LMt|aKTe+Nn<bD~
z^HD+xZrh`BqGAwE13~orTU~wM3=dy=IR|%I<pJEzxPgD=rbV&N-Gp{Bb!tV>is~ac
z^4g`y-E=}I4|Qa*40~-p7QLpozORhKP94#jWnLS-&_PHW9PC+p&&a52g&WJd22&)x
z<89Q-d}I>>#99yL+Ujc-GRu1@D8JB>xU*aB9E|6Qx+$&Y{VV0^$P*22nQMx|))24d
z^jBv%Q{#!2ijh|o)0lI0P=A>R!&GbCBe^Vx)2>GuxMS+?QfP#+UN5CY4AL{$1nKEf
zoM3g)&el2`5t5U<zFb{N9$QLEZ|8fhX@@EBPx)rIa)#b99^=!b2+ZnKjN&5_jBLU7
z*oLTvJ^OGJ?O>AQ=y|BpqUQar=AH#>wwKUrhK{)3T*8>1|4O23ruz?caw)aS4>M(G
z=U})nTWnz!lT>+TM;TXDkfZ8ioTd8mu1D;EYJK6Jq&M@!F_NcUi5GY)Q^b-#P>VKM
z2N{lZo9Q;{LdcE9?xE*h^sB_ap-Vk4xdsKpe^kagN{P>#Bi!6&r^UI&C1D_N%rLX<
ztIE~o)W-Oy)Ks{mow!oj7vU1IkHIvUB(66(GsW2&0z6~E3%SsDlX_`tG4fufbeJip
zx_@}^UiN@<P{5}?ueoA^zxC6x$NJ$TXJ>v*Cx)-=|CQs5Adl}fWwCmTU9jvof^!~O
zqwRT_If*v&h}74@PiF98N<|q;*Xg?EY>Ah**ErWPi%T7VIeUEW6%8Ku@Y_DxdU3Bq
z-EEP$O1YGCa<mFN+p;OFhGSLW%^XIG&LG>}*nqjlB+oAD(=0iewQkycj9HzQ9P)-D
zL%#b-dF6vq*GkMpluSTXTlz$`?*om!N8pkcD)M@TVb+up%Gh5`TWKJmJhN!TytQn|
zH!Z+$K3`X?R_A~~-|Xzo;lwjs|Jeu!QZfEQe3s?m4@%mvN#+@{XvQ<EVr%<?knY_R
zgz=B-gPFfmkw~^T+il@%uUp|NW#TJ!EgpMbkR;qWbu6!}u&J0|I8F+D*m9ohHM{?U
z#S8L%$G#T{hOyD47~1=^VEEq`bjS=_*HTJ`;Fxm&?m~|vioLWu9Ute-)Cy(t0(<65
zdh`{S1NpA_T@9)7<ci(9X8KGQeS_%sC*!WaCREn`ac9}dR@C_V@i+aDD<AZoB4a_o
z>>Y|s#*^<TusTulDfu(gO9Pc$li~b#HZjlTC>)!5qI6RshWuvb_1V?fdT-*Gr8o!}
z7%Q9}bnTGDt7%lu182KpuqXt|gz4AY;&b6st8>nwNUE>zpaF*6iO^#&O%eV5*y%PY
zBteWBT0_^R>-f2m4IMdjy_cGoEx#maxWfv!w!xquAIx$D0rR(rY|r|dS9PUN@S?ZR
zAu7J!!hY<J@FFt>$jD&rD9o!@MZf2yBv>^MsS1uV9?o7})4xIO{zNYb)89V8oF=>6
zZ%T2ZT^~f6r_Ob@=-yngEje)6Cr%}$>Ga8(TvgS>(;315%5XUfk9$vPwTYkec2Dlg
z9Au@3Qi`hjx?I+Gb2RQVmzFaX2ws}mMYRPe66KLrHdD_&{l0^QHXNRu@m#E4XC(nH
z`|h=_d~<K|IuGGLnb>}_oRsEh$LseY&|&^7SSHFO?^D)OF+dLZ4OXMf@e5pkkv*Cz
zXfaz&oYI4=-}#=Izf0VWKATGSf2owadF#9OfPlJJyV>ju?9i&QI&^XM$5i!yZ!P~{
rpOv2*^ncR^{pZh+e=4Yd&dLv|!k^3cXQ2EHl%KivUzb~@!M^_jdD0iu

diff --git a/examples/screenshots/webgpu_depth_texture.jpg b/examples/screenshots/webgpu_depth_texture.jpg
index 35b06935974bb9787ab26daee5a8712bd8be7e9b..001fe203e0b94e476d12e5c1c76aa979a0fbc77f 100644
GIT binary patch
delta 21
dcmZ2ya#d=>CFV^++8mn`ghCiu{xkf)2>@O+2%`W1

literal 6702
zcmeHLX;c$iwyqE$M&JU<prXxHL2zJ@84+<rhK^i7KnQ_W!5AaZT$x7)RLljDDF}!t
zqYNP;B#{sZAR>dvWtKn~0>~T^1QLd%9_U;B=e-~Hd#n3TyK9|&s_L9t^?m2;v+Fwr
zeg=mD87p&3a{vMX0LaDxfc-!eAOhL6O)s)(qR>q*CMpUQ6&Dj1|5jk|E#fdZOk8}6
z#1{D0P1`t=l-Me{dGY-*5iuxKY%5G0_O1D!a$pT0y9HQ<Jb^-F0TEdUR2Bl(L$Wq7
zihYB&0ra~;M4+N#;;;<}5;7aGq0kMmVjD;|dWLVb0iv>E+x8wkCBEJEGHjo(oX(y2
zms|FqF0GZfYv-uwUitkle5-=uj-9(y52&dhJfx>@U}$7~?9AD7=I1Ret?Vy4I67T&
zcDd@|c@2FXbK{nue?VYRaLB!=`wt#QKf)y>CMDxjo~EW{W#{DP<-aN@d{b6VB2(U0
zQ0v~+zi;@^*woz7+11_C`?-%X!WtbL|1vQ-HO*c4y12B=<FBkj0O)r*8~%Ur*&q%P
z*&K7&CLf4M;6{PUii+($D!%QME$p)I_I)~cw#c21e_2`!->+-Ok-zeL`&I=Ny%AOJ
zCZ%tTevHuF{}ZEs5c&rnkO4?SAsZ70l?BWIVdNz6Lge;VK>cU+ZyEx#0xiBCZ`iPA
zDB6r;;iXg`Vah)`VtVU0pTA(^NkZaCKQ<M2q=01Ic>lGDB=@9pZX~A(1m3g>l$_py
z!2151k^F8z`RBd=xG@M*1_79kuq_f;mj(gRbNoJ5rdXIXCnBzZ0RnY~%U<WLIo$*r
zht1W*$C52ZB`X{~jgCLq|MG?OSo4o9av_mE{OA128Sh6R;0bjDf%c}8%K2TuZ;+E>
zfNzN5oGb^4tUwyAM`F;r(}(ll{Zh6YFRd{X*^(uZ9Mfbp>UTH=(NbYP-bliwghz@l
z521<ZpelRIx6Ouho5}!6|17z($TE~*h>0)>lL=m2^WZm)*dHJS@5xKv!IdJ<v^?62
z3t~n42u91DWGv)=w?ZgE9@?QY4YiPlJ}R?ASs^Fd%>SCZqln@u#E^J7Ls>T~SXww`
z>EnA@8Ctb&hq<v2QVX#-5cu3fjbd#NYiuO!KEiia5kD(nB!Gan(FeU08UrOSbeWkf
zp6@4=+(n!^u<Si#dZY@qINy2A<II{>#HW;9`t;~wzF;`NlU@I%LvY;T+U2LG?^@!_
zQc1LOPju~y!yO~>j11y<B%Zo@eQCkd(x~bEWKX89`NXJJorEdoxM3g{`Of!-&OlXp
zz&0#-#^huCdEEw?YSxTgsfE){_`ahfcXV<UH7_ppa;k(gCd*~B^;MHhudJ0y=P)*5
zaJoMpW42oepJ6;s&Yq0i5vCCuNK_%dWH5prFf8@uqgjr39Gaf)9l@b7Gq{{*WFZ9v
z)OaUqI|K$RQ~9jwMQav4#(nF$6f4}UHbRRlOwONX<k+=%=DQ9aYX0-RWM`AE=B?d_
z_t5H_KcBSso~wsy+FBv1_LZ!s1lYx76jK7+eBONB32T_Yr|6De(|SK7(X`xZfHLoF
z>79$qME?Z@bjGzoK+k%-dJ5KTU|><3Ctx-rEDuo0Dd|X}L$?Bb<odb}2t>(Ot8c{U
z`#<R;J1@l$SXy&I5j8R=>KaA^Q*PIk1YX_02U9+(@|@lBWqEBi@SvIoy~u~RlS-e1
zaMVt9W;*tq&<%1NWSGj`m^+dtc&pquk^6aNC2H)-sfA}TyR4n{D2_JG%cM28>A_b?
zof!9U+FWRVM;kjlwrL$ZJmU`Eczd`%i>YI}rH}BrCx;&65}uKy(>k%etmV)tsb?~$
zTDS4`<h~f4_bPa3HYdhW=(j?JOpFFb?Jvdpy+!tNHC{#6u|yJJI#;n0MOT!J_OXsc
z7P}_!(2a}X)FluY^W>fQs8~f=!n(7mYrVeGb5b19^T7xCqr)7kf3cIxy~sKQ_K!=d
z>ih6EgJIEMF4dLR_#D-~FoGXdNynmn(yDq(;%wADTjpyzR1tkNg=vJaM*lX3{c<3u
zi-H&a%AX%zGuE3RqfQ@=+r}tjq-Sb1ewC}fx($Bi5E6JK_jrq5m&&R?r44237ZNO&
z80jxW3H0?&99X~1599<|zwW8tH#2`3TQIGB<xT1l^t0n%c3#^OmK0uL;lV&c(r<{M
z+L60Rm(Py|2+2#~3UiZ*SEEoi4_86pHpM5aGUUKV4?&?B+Ps6HUr?<xsf}cksQAk2
zONeY6Z=OtUN7I8cG*0bMS#;B3iqdORm4#K(fLb<!nATo2%b1P6XMZ=z-^O98N^pF=
z*U!vdkW($s_w>%~&f&`7+H4iHe0_#vU4nAcJ=lX6<6Yi-!jU@_6LcEW&A*^pp2pD<
z=Z9>bS7{W6>KER8+h8BTuJroS<@9Opun$l7Y7Gb|(6)0GZL2$I2Yp{PX=YT;na`i1
zk>bWH-jYmIo~u}zg<_3XB>r9q0>14}oqXwetELXiAs~Qnv(qf0EeTUQe$)P@hWJq7
zFONYWNiOUb2vnd!z%Km2ZL=qben=)kSF1tb-WmG16je=-Bozt*=T(7-%x@V~4)Q_y
z_Rud<Elj;*oyXG#R*PTKoL(gde^K|0O&a)He`%~)yT!r@xBW9`x8=RaUiR@hFg6}(
z__4pT+w?A$VQd^;<UW?<IB2PLB3J7d;m09(m|C$jHcg5ieL2&nD^1v~)UaCqq_|Zy
zfPd7qJ$6*4KDt!r>7UZc<`-O*@+sFyz2~p-zT(2!CEKkqYVCP6LMt|aKTe+Nn<bD~
z^HD+xZrh`BqGAwE13~orTU~wM3=dy=IR|%I<pJEzxPgD=rbV&N-Gp{Bb!tV>is~ac
z^4g`y-E=}I4|Qa*40~-p7QLpozORhKP94#jWnLS-&_PHW9PC+p&&a52g&WJd22&)x
z<89Q-d}I>>#99yL+Ujc-GRu1@D8JB>xU*aB9E|6Qx+$&Y{VV0^$P*22nQMx|))24d
z^jBv%Q{#!2ijh|o)0lI0P=A>R!&GbCBe^Vx)2>GuxMS+?QfP#+UN5CY4AL{$1nKEf
zoM3g)&el2`5t5U<zFb{N9$QLEZ|8fhX@@EBPx)rIa)#b99^=!b2+ZnKjN&5_jBLU7
z*oLTvJ^OGJ?O>AQ=y|BpqUQar=AH#>wwKUrhK{)3T*8>1|4O23ruz?caw)aS4>M(G
z=U})nTWnz!lT>+TM;TXDkfZ8ioTd8mu1D;EYJK6Jq&M@!F_NcUi5GY)Q^b-#P>VKM
z2N{lZo9Q;{LdcE9?xE*h^sB_ap-Vk4xdsKpe^kagN{P>#Bi!6&r^UI&C1D_N%rLX<
ztIE~o)W-Oy)Ks{mow!oj7vU1IkHIvUB(66(GsW2&0z6~E3%SsDlX_`tG4fufbeJip
zx_@}^UiN@<P{5}?ueoA^zxC6x$NJ$TXJ>v*Cx)-=|CQs5Adl}fWwCmTU9jvof^!~O
zqwRT_If*v&h}74@PiF98N<|q;*Xg?EY>Ah**ErWPi%T7VIeUEW6%8Ku@Y_DxdU3Bq
z-EEP$O1YGCa<mFN+p;OFhGSLW%^XIG&LG>}*nqjlB+oAD(=0iewQkycj9HzQ9P)-D
zL%#b-dF6vq*GkMpluSTXTlz$`?*om!N8pkcD)M@TVb+up%Gh5`TWKJmJhN!TytQn|
zH!Z+$K3`X?R_A~~-|Xzo;lwjs|Jeu!QZfEQe3s?m4@%mvN#+@{XvQ<EVr%<?knY_R
zgz=B-gPFfmkw~^T+il@%uUp|NW#TJ!EgpMbkR;qWbu6!}u&J0|I8F+D*m9ohHM{?U
z#S8L%$G#T{hOyD47~1=^VEEq`bjS=_*HTJ`;Fxm&?m~|vioLWu9Ute-)Cy(t0(<65
zdh`{S1NpA_T@9)7<ci(9X8KGQeS_%sC*!WaCREn`ac9}dR@C_V@i+aDD<AZoB4a_o
z>>Y|s#*^<TusTulDfu(gO9Pc$li~b#HZjlTC>)!5qI6RshWuvb_1V?fdT-*Gr8o!}
z7%Q9}bnTGDt7%lu182KpuqXt|gz4AY;&b6st8>nwNUE>zpaF*6iO^#&O%eV5*y%PY
zBteWBT0_^R>-f2m4IMdjy_cGoEx#maxWfv!w!xquAIx$D0rR(rY|r|dS9PUN@S?ZR
zAu7J!!hY<J@FFt>$jD&rD9o!@MZf2yBv>^MsS1uV9?o7})4xIO{zNYb)89V8oF=>6
zZ%T2ZT^~f6r_Ob@=-yngEje)6Cr%}$>Ga8(TvgS>(;315%5XUfk9$vPwTYkec2Dlg
z9Au@3Qi`hjx?I+Gb2RQVmzFaX2ws}mMYRPe66KLrHdD_&{l0^QHXNRu@m#E4XC(nH
z`|h=_d~<K|IuGGLnb>}_oRsEh$LseY&|&^7SSHFO?^D)OF+dLZ4OXMf@e5pkkv*Cz
zXfaz&oYI4=-}#=Izf0VWKATGSf2owadF#9OfPlJJyV>ju?9i&QI&^XM$5i!yZ!P~{
rpOv2*^ncR^{pZh+e=4Yd&dLv|!k^3cXQ2EHl%KivUzb~@!M^_jdD0iu

diff --git a/examples/screenshots/webgpu_instance_mesh.jpg b/examples/screenshots/webgpu_instance_mesh.jpg
index 35b06935974bb9787ab26daee5a8712bd8be7e9b..001fe203e0b94e476d12e5c1c76aa979a0fbc77f 100644
GIT binary patch
delta 21
dcmZ2ya#d=>CFV^++8mn`ghCiu{xkf)2>@O+2%`W1

literal 6702
zcmeHLX;c$iwyqE$M&JU<prXxHL2zJ@84+<rhK^i7KnQ_W!5AaZT$x7)RLljDDF}!t
zqYNP;B#{sZAR>dvWtKn~0>~T^1QLd%9_U;B=e-~Hd#n3TyK9|&s_L9t^?m2;v+Fwr
zeg=mD87p&3a{vMX0LaDxfc-!eAOhL6O)s)(qR>q*CMpUQ6&Dj1|5jk|E#fdZOk8}6
z#1{D0P1`t=l-Me{dGY-*5iuxKY%5G0_O1D!a$pT0y9HQ<Jb^-F0TEdUR2Bl(L$Wq7
zihYB&0ra~;M4+N#;;;<}5;7aGq0kMmVjD;|dWLVb0iv>E+x8wkCBEJEGHjo(oX(y2
zms|FqF0GZfYv-uwUitkle5-=uj-9(y52&dhJfx>@U}$7~?9AD7=I1Ret?Vy4I67T&
zcDd@|c@2FXbK{nue?VYRaLB!=`wt#QKf)y>CMDxjo~EW{W#{DP<-aN@d{b6VB2(U0
zQ0v~+zi;@^*woz7+11_C`?-%X!WtbL|1vQ-HO*c4y12B=<FBkj0O)r*8~%Ur*&q%P
z*&K7&CLf4M;6{PUii+($D!%QME$p)I_I)~cw#c21e_2`!->+-Ok-zeL`&I=Ny%AOJ
zCZ%tTevHuF{}ZEs5c&rnkO4?SAsZ70l?BWIVdNz6Lge;VK>cU+ZyEx#0xiBCZ`iPA
zDB6r;;iXg`Vah)`VtVU0pTA(^NkZaCKQ<M2q=01Ic>lGDB=@9pZX~A(1m3g>l$_py
z!2151k^F8z`RBd=xG@M*1_79kuq_f;mj(gRbNoJ5rdXIXCnBzZ0RnY~%U<WLIo$*r
zht1W*$C52ZB`X{~jgCLq|MG?OSo4o9av_mE{OA128Sh6R;0bjDf%c}8%K2TuZ;+E>
zfNzN5oGb^4tUwyAM`F;r(}(ll{Zh6YFRd{X*^(uZ9Mfbp>UTH=(NbYP-bliwghz@l
z521<ZpelRIx6Ouho5}!6|17z($TE~*h>0)>lL=m2^WZm)*dHJS@5xKv!IdJ<v^?62
z3t~n42u91DWGv)=w?ZgE9@?QY4YiPlJ}R?ASs^Fd%>SCZqln@u#E^J7Ls>T~SXww`
z>EnA@8Ctb&hq<v2QVX#-5cu3fjbd#NYiuO!KEiia5kD(nB!Gan(FeU08UrOSbeWkf
zp6@4=+(n!^u<Si#dZY@qINy2A<II{>#HW;9`t;~wzF;`NlU@I%LvY;T+U2LG?^@!_
zQc1LOPju~y!yO~>j11y<B%Zo@eQCkd(x~bEWKX89`NXJJorEdoxM3g{`Of!-&OlXp
zz&0#-#^huCdEEw?YSxTgsfE){_`ahfcXV<UH7_ppa;k(gCd*~B^;MHhudJ0y=P)*5
zaJoMpW42oepJ6;s&Yq0i5vCCuNK_%dWH5prFf8@uqgjr39Gaf)9l@b7Gq{{*WFZ9v
z)OaUqI|K$RQ~9jwMQav4#(nF$6f4}UHbRRlOwONX<k+=%=DQ9aYX0-RWM`AE=B?d_
z_t5H_KcBSso~wsy+FBv1_LZ!s1lYx76jK7+eBONB32T_Yr|6De(|SK7(X`xZfHLoF
z>79$qME?Z@bjGzoK+k%-dJ5KTU|><3Ctx-rEDuo0Dd|X}L$?Bb<odb}2t>(Ot8c{U
z`#<R;J1@l$SXy&I5j8R=>KaA^Q*PIk1YX_02U9+(@|@lBWqEBi@SvIoy~u~RlS-e1
zaMVt9W;*tq&<%1NWSGj`m^+dtc&pquk^6aNC2H)-sfA}TyR4n{D2_JG%cM28>A_b?
zof!9U+FWRVM;kjlwrL$ZJmU`Eczd`%i>YI}rH}BrCx;&65}uKy(>k%etmV)tsb?~$
zTDS4`<h~f4_bPa3HYdhW=(j?JOpFFb?Jvdpy+!tNHC{#6u|yJJI#;n0MOT!J_OXsc
z7P}_!(2a}X)FluY^W>fQs8~f=!n(7mYrVeGb5b19^T7xCqr)7kf3cIxy~sKQ_K!=d
z>ih6EgJIEMF4dLR_#D-~FoGXdNynmn(yDq(;%wADTjpyzR1tkNg=vJaM*lX3{c<3u
zi-H&a%AX%zGuE3RqfQ@=+r}tjq-Sb1ewC}fx($Bi5E6JK_jrq5m&&R?r44237ZNO&
z80jxW3H0?&99X~1599<|zwW8tH#2`3TQIGB<xT1l^t0n%c3#^OmK0uL;lV&c(r<{M
z+L60Rm(Py|2+2#~3UiZ*SEEoi4_86pHpM5aGUUKV4?&?B+Ps6HUr?<xsf}cksQAk2
zONeY6Z=OtUN7I8cG*0bMS#;B3iqdORm4#K(fLb<!nATo2%b1P6XMZ=z-^O98N^pF=
z*U!vdkW($s_w>%~&f&`7+H4iHe0_#vU4nAcJ=lX6<6Yi-!jU@_6LcEW&A*^pp2pD<
z=Z9>bS7{W6>KER8+h8BTuJroS<@9Opun$l7Y7Gb|(6)0GZL2$I2Yp{PX=YT;na`i1
zk>bWH-jYmIo~u}zg<_3XB>r9q0>14}oqXwetELXiAs~Qnv(qf0EeTUQe$)P@hWJq7
zFONYWNiOUb2vnd!z%Km2ZL=qben=)kSF1tb-WmG16je=-Bozt*=T(7-%x@V~4)Q_y
z_Rud<Elj;*oyXG#R*PTKoL(gde^K|0O&a)He`%~)yT!r@xBW9`x8=RaUiR@hFg6}(
z__4pT+w?A$VQd^;<UW?<IB2PLB3J7d;m09(m|C$jHcg5ieL2&nD^1v~)UaCqq_|Zy
zfPd7qJ$6*4KDt!r>7UZc<`-O*@+sFyz2~p-zT(2!CEKkqYVCP6LMt|aKTe+Nn<bD~
z^HD+xZrh`BqGAwE13~orTU~wM3=dy=IR|%I<pJEzxPgD=rbV&N-Gp{Bb!tV>is~ac
z^4g`y-E=}I4|Qa*40~-p7QLpozORhKP94#jWnLS-&_PHW9PC+p&&a52g&WJd22&)x
z<89Q-d}I>>#99yL+Ujc-GRu1@D8JB>xU*aB9E|6Qx+$&Y{VV0^$P*22nQMx|))24d
z^jBv%Q{#!2ijh|o)0lI0P=A>R!&GbCBe^Vx)2>GuxMS+?QfP#+UN5CY4AL{$1nKEf
zoM3g)&el2`5t5U<zFb{N9$QLEZ|8fhX@@EBPx)rIa)#b99^=!b2+ZnKjN&5_jBLU7
z*oLTvJ^OGJ?O>AQ=y|BpqUQar=AH#>wwKUrhK{)3T*8>1|4O23ruz?caw)aS4>M(G
z=U})nTWnz!lT>+TM;TXDkfZ8ioTd8mu1D;EYJK6Jq&M@!F_NcUi5GY)Q^b-#P>VKM
z2N{lZo9Q;{LdcE9?xE*h^sB_ap-Vk4xdsKpe^kagN{P>#Bi!6&r^UI&C1D_N%rLX<
ztIE~o)W-Oy)Ks{mow!oj7vU1IkHIvUB(66(GsW2&0z6~E3%SsDlX_`tG4fufbeJip
zx_@}^UiN@<P{5}?ueoA^zxC6x$NJ$TXJ>v*Cx)-=|CQs5Adl}fWwCmTU9jvof^!~O
zqwRT_If*v&h}74@PiF98N<|q;*Xg?EY>Ah**ErWPi%T7VIeUEW6%8Ku@Y_DxdU3Bq
z-EEP$O1YGCa<mFN+p;OFhGSLW%^XIG&LG>}*nqjlB+oAD(=0iewQkycj9HzQ9P)-D
zL%#b-dF6vq*GkMpluSTXTlz$`?*om!N8pkcD)M@TVb+up%Gh5`TWKJmJhN!TytQn|
zH!Z+$K3`X?R_A~~-|Xzo;lwjs|Jeu!QZfEQe3s?m4@%mvN#+@{XvQ<EVr%<?knY_R
zgz=B-gPFfmkw~^T+il@%uUp|NW#TJ!EgpMbkR;qWbu6!}u&J0|I8F+D*m9ohHM{?U
z#S8L%$G#T{hOyD47~1=^VEEq`bjS=_*HTJ`;Fxm&?m~|vioLWu9Ute-)Cy(t0(<65
zdh`{S1NpA_T@9)7<ci(9X8KGQeS_%sC*!WaCREn`ac9}dR@C_V@i+aDD<AZoB4a_o
z>>Y|s#*^<TusTulDfu(gO9Pc$li~b#HZjlTC>)!5qI6RshWuvb_1V?fdT-*Gr8o!}
z7%Q9}bnTGDt7%lu182KpuqXt|gz4AY;&b6st8>nwNUE>zpaF*6iO^#&O%eV5*y%PY
zBteWBT0_^R>-f2m4IMdjy_cGoEx#maxWfv!w!xquAIx$D0rR(rY|r|dS9PUN@S?ZR
zAu7J!!hY<J@FFt>$jD&rD9o!@MZf2yBv>^MsS1uV9?o7})4xIO{zNYb)89V8oF=>6
zZ%T2ZT^~f6r_Ob@=-yngEje)6Cr%}$>Ga8(TvgS>(;315%5XUfk9$vPwTYkec2Dlg
z9Au@3Qi`hjx?I+Gb2RQVmzFaX2ws}mMYRPe66KLrHdD_&{l0^QHXNRu@m#E4XC(nH
z`|h=_d~<K|IuGGLnb>}_oRsEh$LseY&|&^7SSHFO?^D)OF+dLZ4OXMf@e5pkkv*Cz
zXfaz&oYI4=-}#=Izf0VWKATGSf2owadF#9OfPlJJyV>ju?9i&QI&^XM$5i!yZ!P~{
rpOv2*^ncR^{pZh+e=4Yd&dLv|!k^3cXQ2EHl%KivUzb~@!M^_jdD0iu

diff --git a/examples/screenshots/webgpu_instance_uniform.jpg b/examples/screenshots/webgpu_instance_uniform.jpg
index ed64a2bef3e0f09dbe137e4937c52ffbc222a0c0..001fe203e0b94e476d12e5c1c76aa979a0fbc77f 100644
GIT binary patch
delta 51
fcmZoLy(+cg67%K>OiqmWiOCa~HnRL@_<s`s+F%ls

literal 6784
zcmeHKX;c&0wk`+|CL2K!0Ra(Yl0lK7QGr$z2wY?ajRaICF$9!(j*6g+$`lkunn^?i
znju7HWi&yh0hxyoXb=gI1PlofQa9+m{pYRqUibTP|6F&ivv*aUQ)_+SuD$o!2OI&X
zf!!9SFjIhw3jnw{7XXd}cL8p$9of;`JHi9m(Y!o75FS2WKE7`QD!|VN6@c>b@eA<_
z2<`~yPFP4#c<15!XWYCH2(KWN5BhETZz=E%Aj%JHay^7_i2~fBTo6$%u!F0BgOT?e
zG!D@3!o>~Y;pKyJAPDW|z=l9LV0k%6IV+<%a{!MhubBL?(|mhuuRsso+^c;%EuUZE
zOx0U)yFRMo@v8xM1O)d<?3a`}tfZ`>s-~l>r~k9T36rzuOwG(;7WNK~PR<uD!reVQ
zy}YmaTn`Ki4hanlkBGe+7k}@5LSlMGW)>#<(c_$g!lL5mB`-=}R@c;G>u|5@8`|1C
zI=i~x;R*c%p9Y79M@Gj;GqZop&3|6_vPfH9`$}JDFgG^20LXVb9Q{A|aENnp@Aw?L
z!-tDIlrtcrJiPM9_{2`zLa*H1b4dF(|K2lc`BiTP6pq_b#jghR3GP$WAswdeQ2NH`
zuMxWQ|6=qfp+EV6V}LM(i}S-EqJS~LMlS%ra^DgJRDPKMkPD1UwU|1LY5h0E`^Hpr
zlypaoA@dl?Fwn+d7MX@+SCYn&j}z5Pv6k=Remf~l`$eu8O~r#ib?>&cb2|vI6nxRl
zK|t=uyZ<^LM9F~wRGZz423R5>z;li{I+Mp6B|?pPP&x(zZTjmdGfV1Vc@uSorin?Z
zGoKZ%ck=xCRJ=m|uOfdC{#uf=D_f_3r)-$`C4hh@#1#bk@W05Fd;)B^e&GcIB8c5u
zaWVU88lUg7tY@$TnqRlP>YFp{YICN>gl<$-<yI)3=1<Zt5YS)Jh9-+GEEYxoG4=%n
zAgsYgGAcJSCkeJf-$1g@Z`ZbVN;gnuac_4STW2`8GVvsH4f&camtp&`_)&Mu=(^s|
z`?PL7@z;>PIV)lFbmnYN`i&4wPVt%_J_1*OaD5W%qeQf13sP3#E$?b(j1xg%SCqo^
zmfQUp1zO7!AJSG=?7Lx)dV3Ei-*7w0F82!)S=^@b0-lmR!ljWJHkg!2<a2y0-WY)?
z-xb*x054{Tu%?=E5lVLTOc0<RBA`sC)@I%AUd=Fz7r+K(cVJ%nx<4|#5`H^_VKH+p
z5%cDXn)}IkxK$S!IsSQk*o*-^yG355F!LF$vPMv!2JiKuLfVry%zBIqWxbwt&NmGe
zAB201lju@-W<KMdmo@1?dL)?eaqqZY2o9aMVO+oNo*0;|9*m|E=UeC#UL`aTm^Mm6
zqN86$`BGDgq+Eumi-9i8^dX(vbTnN+)dxr^dOIS`Th(=aPSPwG1>Yziep%D7pb|q<
z-evDRP?xiLz2&PbGl6!A)Y|3-?U|({$m$N%M$0N{tdo0gm8p(=6gxKLtoF)r_i;TV
z8><7AV=BRZL*yIqPP}gxWp3gQW7hv|ik{_*V7a80{kTXCD;;vj6Sq1O-dVqK-|1Go
z>IgHLlseGOmuVdChwu+7A1oj4d0}{%!02N=>-{`XUrYi48Kzuk&a>izFy}6~Q%Pi@
z3%_oXm9cH2kFV$9Td5wVN=gJA;f-A8rc_82Zb8J%%VUHU#*}#p?~Z=X@}d@Z9KTK+
zy=eu#U!Lo__28?G$@GVJ+anwG2khXli`TKPNUO$`x|PPe%^&?x8m-xtsJscEFzFfP
zTQSPlC^K^RjCA_x<NOsxb(l3+OX`eH&hdN6)nX9UE{9cZ!)x%i<CkDcc|mk5L2`~Q
zdEdY)GRPaw3Sufmt(m#g9<<fyHj^(-sJqJ7xr-B;RDQ3Mn|o;3*^qfE_LPZaFd-oM
zYSB~*!iemQwp3;qy$`9ak67OrUmF-cqam>#%Y@FHEIT4666H^gV{K-3$()eZ5*s5h
z?1ViPgyc2b732fLOSEo|!F|w}m+WlomFNV!BxmdLmIxp6%YhG}V#Z_A)57=?m`Jx4
zv|*8Q;){BV67~>1e@tmNdyNsN|9rZByg4vAF~moI6tQ`P{H1W~afx~*G8os3t4#9}
zNhS(W*n4|v?IyI`=j9~JXid~PZ0>ICSAUdPjX5pp#<HqSvfe>2LRME;MUyN<Jhj`t
z&GrFB4Q8r}-8+$+-K<oQ;UChEKqR!VO<nW0tfKfneGLyDw@!#ar%54TVgy+OzggUa
z4>^~v<(|#9eHHGnEgp1PyJ)z~{w^~yE7;<K4`Q_hsVkx1BW|EHE1+MZJ^_y`x?0vB
zMF4@HKp+6E?^TNkBA)rc=mvqfeIc>-wX_49#o6nq&aLaNa{ZdvBWLf3XBBxp>sM8<
z-%4<hzx=VM%H$X@XC1%xxObG0b9bZgCwNAjK_IMCvCz&o!}aSWpNq=x$6(ci<`x#m
z23t54vOTK39WC1;&8%^$qt2)3)%vYWD8hnChNhE__v=q0=5UVpK_GvE`#^;ta03C^
zUm^41Xjr)O%lsahWv9&63{Od%){KeA$`hnHp}))?rP*MZcWHdUn`!iNJ1fj9pHm|D
za#5iy85s~zQDJ`OFzaO(ZCU~7gAU&-=-&U`K<UvSz;_Q#kPfomQjub?Q!RjPqwftG
z<PlAxO$5R$L$uN~t7Oa9n=LJ3ts5e(MTwVwjhD;`OZM4INVBLue)KVS&qw{)L??zz
zWG$OA&@f7uM!hB$FN1(T90bybqjio3`;R(LcVQJ*Cr85W=+*awK+Vdoyj0!DXexU)
zP)_X&I)dnR>lBgfq9i%8uLlm(?dz=6yc96%Tkh#mtAy|DT%w?jJQXRESmiQGbU)5|
zRwTYff<}z{l^S|K&rvUA?X<eJ=ZsK~<t^&NzN*)E93J04C}#OIEhFP>^OsV(^SIT3
zwZGC}om9W3p<Hx`4}HxN@g9jDoh8LRmC9{SK0$9-?7hw~_NHM<B@@a?=7*30xPxQ3
zEU7}BlfsGo0iy=8WuMW#)L~Rp?dd)(#TM9P#A0P=`Fokjb>;2HO_r$Ai)6W<Jkn(l
z3RY)S=w#-QgBgsQq|D`jKj8QFhD_f@oM?`P>bc^Y(2Xrlc>J*(aS^q@$k4vk(Lbx8
zoN+FCl^yQ-$*o#sY1{h+GN&$WnmnmZ?#L;AR#-l9jji}e{n`BdkcqnX9-akfv10)e
zf=dCU6c>-#BNfZ6NeW}~OI<-Vwy8eG$#T^a@v(Box?Kdyf>UW!T*~E7vhQoebSC{)
z+Uy(jQ?fr`F$Q|oCsh&S=BHwB+oJqeH+(jWw(X*!O)jpzo9w2FI^Vfd7q{eNdZ$i3
zAkZpfHR+CstA=)R%OawSVoxF;u_nt!eR|XDpV_%vz39ZNneQ)JR>Aw#t-^i2tdwbF
z)&4?{jU5F6XL__}V`F3J3FCWtd@To9M+Q%wE@w)cVEQSYCu$ZWme#I2OY8~^=tu{F
zmb7(|N4FBN4wpGcAC(Go6^(2<bKG@nMPw@b_(V_Ax-Zk3w|88jW;m*C!NZG|`XXx3
z$!qng^i!D7zLPV)T`Yg6vDZPw>STABtn$_5t2p-3p%&-KI{aZ(3+vEEGR?wG!&|uo
zeb_I%O%_T8f!q1bb<~CC3lzF!Wl18%@pXd(#%W?kUMKb71A{{<>s>`pRq&;j7{Mc+
z9?fDft+VN{AR$e!t;Ow2jWOkoxZ%?FFUD1rH~u1UGbE`QdSM_r;%m<D&jvcj9;RC|
z(x*&wob;i~03PW#ZxGt`kIR_`zA3IxFW1^8C61q^e<e0=`Pye6XJDIKg_;@UKDK$u
zbnnOjsU+4f%HSvsos2bS=|9(f<<x*nn8zI15WsdYEE=Qw3;K8eM)Q4i@(Ps4aag?&
zrI;QaGCzzMpAYq;AmMER89qb@N=)O_@IuBA0Z)$1HI1bTk-FpQlUf}OHpA@kqUe&P
z3>Rj%{nMXwuU_Qo(P-Cr*ReZJ?4H0}#fR4g%XL#YCl+Tf4JG=$ixb};pdmi6m6c7H
z)FU=(`MB}3b`?hHzh=Tv-e}F+!S1?GhqA}!Cg10IX9?<9>h#yXXoNWmP9J=OcrD%+
z-#K$MEGI1S{J`^11xwGWLaio^Q0N>6`6gafkB9&Pt?ZuRkbr@huGSjg>M<APcF#LV
zY1Emi{zxnMisYe0`g3HQ$btOdtUAWSx2Hg0a_LQXUw3L)v8wX@hK%Q9H)E{^Y;8hr
zBF~!>H<l@D*WJ~11uDbcVk$O^u#wGT8f3TX$~P3T9@NblW8p=qyXzR79@S-vnQeo)
zT+^;2W{_5*-oIf_NIJ6T(2*;WXDW{6cO5t#XXl`1DtKO__<XwSQla9jd{5bHUFTH;
z@MZXujSKnv@i6}&HqyVb#P3OsdmE=ij*(=lk<A;ZmG|bVr1m|+NO?9%WxcucvyHh#
zkg2uiMF8VA)li$)<o{$|w>`Oc%_T~$Th!Ka4FA!p>(qxv*Z#TrM27oH=Zl5fzCD6E
zr9}nh{$x5ja^+)1Q=rYavdo0Ae@X>`%)L>8AW-iO0(Q|#w~QaE{iOpnWU~bXVoh4-
zb~Pw(XYPsw0kgwEOy0L%`~TP}cG~f8OXL9b+dkszod>CChB^o=S{MBeUFH4v?7tTM
iFh3mfXT<ynh9AN3BN+Zxk^PY}KT_tuFlD&F(SHJoqEoE^

diff --git a/examples/screenshots/webgpu_lights_custom.jpg b/examples/screenshots/webgpu_lights_custom.jpg
index 35b06935974bb9787ab26daee5a8712bd8be7e9b..001fe203e0b94e476d12e5c1c76aa979a0fbc77f 100644
GIT binary patch
delta 21
dcmZ2ya#d=>CFV^++8mn`ghCiu{xkf)2>@O+2%`W1

literal 6702
zcmeHLX;c$iwyqE$M&JU<prXxHL2zJ@84+<rhK^i7KnQ_W!5AaZT$x7)RLljDDF}!t
zqYNP;B#{sZAR>dvWtKn~0>~T^1QLd%9_U;B=e-~Hd#n3TyK9|&s_L9t^?m2;v+Fwr
zeg=mD87p&3a{vMX0LaDxfc-!eAOhL6O)s)(qR>q*CMpUQ6&Dj1|5jk|E#fdZOk8}6
z#1{D0P1`t=l-Me{dGY-*5iuxKY%5G0_O1D!a$pT0y9HQ<Jb^-F0TEdUR2Bl(L$Wq7
zihYB&0ra~;M4+N#;;;<}5;7aGq0kMmVjD;|dWLVb0iv>E+x8wkCBEJEGHjo(oX(y2
zms|FqF0GZfYv-uwUitkle5-=uj-9(y52&dhJfx>@U}$7~?9AD7=I1Ret?Vy4I67T&
zcDd@|c@2FXbK{nue?VYRaLB!=`wt#QKf)y>CMDxjo~EW{W#{DP<-aN@d{b6VB2(U0
zQ0v~+zi;@^*woz7+11_C`?-%X!WtbL|1vQ-HO*c4y12B=<FBkj0O)r*8~%Ur*&q%P
z*&K7&CLf4M;6{PUii+($D!%QME$p)I_I)~cw#c21e_2`!->+-Ok-zeL`&I=Ny%AOJ
zCZ%tTevHuF{}ZEs5c&rnkO4?SAsZ70l?BWIVdNz6Lge;VK>cU+ZyEx#0xiBCZ`iPA
zDB6r;;iXg`Vah)`VtVU0pTA(^NkZaCKQ<M2q=01Ic>lGDB=@9pZX~A(1m3g>l$_py
z!2151k^F8z`RBd=xG@M*1_79kuq_f;mj(gRbNoJ5rdXIXCnBzZ0RnY~%U<WLIo$*r
zht1W*$C52ZB`X{~jgCLq|MG?OSo4o9av_mE{OA128Sh6R;0bjDf%c}8%K2TuZ;+E>
zfNzN5oGb^4tUwyAM`F;r(}(ll{Zh6YFRd{X*^(uZ9Mfbp>UTH=(NbYP-bliwghz@l
z521<ZpelRIx6Ouho5}!6|17z($TE~*h>0)>lL=m2^WZm)*dHJS@5xKv!IdJ<v^?62
z3t~n42u91DWGv)=w?ZgE9@?QY4YiPlJ}R?ASs^Fd%>SCZqln@u#E^J7Ls>T~SXww`
z>EnA@8Ctb&hq<v2QVX#-5cu3fjbd#NYiuO!KEiia5kD(nB!Gan(FeU08UrOSbeWkf
zp6@4=+(n!^u<Si#dZY@qINy2A<II{>#HW;9`t;~wzF;`NlU@I%LvY;T+U2LG?^@!_
zQc1LOPju~y!yO~>j11y<B%Zo@eQCkd(x~bEWKX89`NXJJorEdoxM3g{`Of!-&OlXp
zz&0#-#^huCdEEw?YSxTgsfE){_`ahfcXV<UH7_ppa;k(gCd*~B^;MHhudJ0y=P)*5
zaJoMpW42oepJ6;s&Yq0i5vCCuNK_%dWH5prFf8@uqgjr39Gaf)9l@b7Gq{{*WFZ9v
z)OaUqI|K$RQ~9jwMQav4#(nF$6f4}UHbRRlOwONX<k+=%=DQ9aYX0-RWM`AE=B?d_
z_t5H_KcBSso~wsy+FBv1_LZ!s1lYx76jK7+eBONB32T_Yr|6De(|SK7(X`xZfHLoF
z>79$qME?Z@bjGzoK+k%-dJ5KTU|><3Ctx-rEDuo0Dd|X}L$?Bb<odb}2t>(Ot8c{U
z`#<R;J1@l$SXy&I5j8R=>KaA^Q*PIk1YX_02U9+(@|@lBWqEBi@SvIoy~u~RlS-e1
zaMVt9W;*tq&<%1NWSGj`m^+dtc&pquk^6aNC2H)-sfA}TyR4n{D2_JG%cM28>A_b?
zof!9U+FWRVM;kjlwrL$ZJmU`Eczd`%i>YI}rH}BrCx;&65}uKy(>k%etmV)tsb?~$
zTDS4`<h~f4_bPa3HYdhW=(j?JOpFFb?Jvdpy+!tNHC{#6u|yJJI#;n0MOT!J_OXsc
z7P}_!(2a}X)FluY^W>fQs8~f=!n(7mYrVeGb5b19^T7xCqr)7kf3cIxy~sKQ_K!=d
z>ih6EgJIEMF4dLR_#D-~FoGXdNynmn(yDq(;%wADTjpyzR1tkNg=vJaM*lX3{c<3u
zi-H&a%AX%zGuE3RqfQ@=+r}tjq-Sb1ewC}fx($Bi5E6JK_jrq5m&&R?r44237ZNO&
z80jxW3H0?&99X~1599<|zwW8tH#2`3TQIGB<xT1l^t0n%c3#^OmK0uL;lV&c(r<{M
z+L60Rm(Py|2+2#~3UiZ*SEEoi4_86pHpM5aGUUKV4?&?B+Ps6HUr?<xsf}cksQAk2
zONeY6Z=OtUN7I8cG*0bMS#;B3iqdORm4#K(fLb<!nATo2%b1P6XMZ=z-^O98N^pF=
z*U!vdkW($s_w>%~&f&`7+H4iHe0_#vU4nAcJ=lX6<6Yi-!jU@_6LcEW&A*^pp2pD<
z=Z9>bS7{W6>KER8+h8BTuJroS<@9Opun$l7Y7Gb|(6)0GZL2$I2Yp{PX=YT;na`i1
zk>bWH-jYmIo~u}zg<_3XB>r9q0>14}oqXwetELXiAs~Qnv(qf0EeTUQe$)P@hWJq7
zFONYWNiOUb2vnd!z%Km2ZL=qben=)kSF1tb-WmG16je=-Bozt*=T(7-%x@V~4)Q_y
z_Rud<Elj;*oyXG#R*PTKoL(gde^K|0O&a)He`%~)yT!r@xBW9`x8=RaUiR@hFg6}(
z__4pT+w?A$VQd^;<UW?<IB2PLB3J7d;m09(m|C$jHcg5ieL2&nD^1v~)UaCqq_|Zy
zfPd7qJ$6*4KDt!r>7UZc<`-O*@+sFyz2~p-zT(2!CEKkqYVCP6LMt|aKTe+Nn<bD~
z^HD+xZrh`BqGAwE13~orTU~wM3=dy=IR|%I<pJEzxPgD=rbV&N-Gp{Bb!tV>is~ac
z^4g`y-E=}I4|Qa*40~-p7QLpozORhKP94#jWnLS-&_PHW9PC+p&&a52g&WJd22&)x
z<89Q-d}I>>#99yL+Ujc-GRu1@D8JB>xU*aB9E|6Qx+$&Y{VV0^$P*22nQMx|))24d
z^jBv%Q{#!2ijh|o)0lI0P=A>R!&GbCBe^Vx)2>GuxMS+?QfP#+UN5CY4AL{$1nKEf
zoM3g)&el2`5t5U<zFb{N9$QLEZ|8fhX@@EBPx)rIa)#b99^=!b2+ZnKjN&5_jBLU7
z*oLTvJ^OGJ?O>AQ=y|BpqUQar=AH#>wwKUrhK{)3T*8>1|4O23ruz?caw)aS4>M(G
z=U})nTWnz!lT>+TM;TXDkfZ8ioTd8mu1D;EYJK6Jq&M@!F_NcUi5GY)Q^b-#P>VKM
z2N{lZo9Q;{LdcE9?xE*h^sB_ap-Vk4xdsKpe^kagN{P>#Bi!6&r^UI&C1D_N%rLX<
ztIE~o)W-Oy)Ks{mow!oj7vU1IkHIvUB(66(GsW2&0z6~E3%SsDlX_`tG4fufbeJip
zx_@}^UiN@<P{5}?ueoA^zxC6x$NJ$TXJ>v*Cx)-=|CQs5Adl}fWwCmTU9jvof^!~O
zqwRT_If*v&h}74@PiF98N<|q;*Xg?EY>Ah**ErWPi%T7VIeUEW6%8Ku@Y_DxdU3Bq
z-EEP$O1YGCa<mFN+p;OFhGSLW%^XIG&LG>}*nqjlB+oAD(=0iewQkycj9HzQ9P)-D
zL%#b-dF6vq*GkMpluSTXTlz$`?*om!N8pkcD)M@TVb+up%Gh5`TWKJmJhN!TytQn|
zH!Z+$K3`X?R_A~~-|Xzo;lwjs|Jeu!QZfEQe3s?m4@%mvN#+@{XvQ<EVr%<?knY_R
zgz=B-gPFfmkw~^T+il@%uUp|NW#TJ!EgpMbkR;qWbu6!}u&J0|I8F+D*m9ohHM{?U
z#S8L%$G#T{hOyD47~1=^VEEq`bjS=_*HTJ`;Fxm&?m~|vioLWu9Ute-)Cy(t0(<65
zdh`{S1NpA_T@9)7<ci(9X8KGQeS_%sC*!WaCREn`ac9}dR@C_V@i+aDD<AZoB4a_o
z>>Y|s#*^<TusTulDfu(gO9Pc$li~b#HZjlTC>)!5qI6RshWuvb_1V?fdT-*Gr8o!}
z7%Q9}bnTGDt7%lu182KpuqXt|gz4AY;&b6st8>nwNUE>zpaF*6iO^#&O%eV5*y%PY
zBteWBT0_^R>-f2m4IMdjy_cGoEx#maxWfv!w!xquAIx$D0rR(rY|r|dS9PUN@S?ZR
zAu7J!!hY<J@FFt>$jD&rD9o!@MZf2yBv>^MsS1uV9?o7})4xIO{zNYb)89V8oF=>6
zZ%T2ZT^~f6r_Ob@=-yngEje)6Cr%}$>Ga8(TvgS>(;315%5XUfk9$vPwTYkec2Dlg
z9Au@3Qi`hjx?I+Gb2RQVmzFaX2ws}mMYRPe66KLrHdD_&{l0^QHXNRu@m#E4XC(nH
z`|h=_d~<K|IuGGLnb>}_oRsEh$LseY&|&^7SSHFO?^D)OF+dLZ4OXMf@e5pkkv*Cz
zXfaz&oYI4=-}#=Izf0VWKATGSf2owadF#9OfPlJJyV>ju?9i&QI&^XM$5i!yZ!P~{
rpOv2*^ncR^{pZh+e=4Yd&dLv|!k^3cXQ2EHl%KivUzb~@!M^_jdD0iu

diff --git a/examples/screenshots/webgpu_lights_selective.jpg b/examples/screenshots/webgpu_lights_selective.jpg
index 35b06935974bb9787ab26daee5a8712bd8be7e9b..001fe203e0b94e476d12e5c1c76aa979a0fbc77f 100644
GIT binary patch
delta 21
dcmZ2ya#d=>CFV^++8mn`ghCiu{xkf)2>@O+2%`W1

literal 6702
zcmeHLX;c$iwyqE$M&JU<prXxHL2zJ@84+<rhK^i7KnQ_W!5AaZT$x7)RLljDDF}!t
zqYNP;B#{sZAR>dvWtKn~0>~T^1QLd%9_U;B=e-~Hd#n3TyK9|&s_L9t^?m2;v+Fwr
zeg=mD87p&3a{vMX0LaDxfc-!eAOhL6O)s)(qR>q*CMpUQ6&Dj1|5jk|E#fdZOk8}6
z#1{D0P1`t=l-Me{dGY-*5iuxKY%5G0_O1D!a$pT0y9HQ<Jb^-F0TEdUR2Bl(L$Wq7
zihYB&0ra~;M4+N#;;;<}5;7aGq0kMmVjD;|dWLVb0iv>E+x8wkCBEJEGHjo(oX(y2
zms|FqF0GZfYv-uwUitkle5-=uj-9(y52&dhJfx>@U}$7~?9AD7=I1Ret?Vy4I67T&
zcDd@|c@2FXbK{nue?VYRaLB!=`wt#QKf)y>CMDxjo~EW{W#{DP<-aN@d{b6VB2(U0
zQ0v~+zi;@^*woz7+11_C`?-%X!WtbL|1vQ-HO*c4y12B=<FBkj0O)r*8~%Ur*&q%P
z*&K7&CLf4M;6{PUii+($D!%QME$p)I_I)~cw#c21e_2`!->+-Ok-zeL`&I=Ny%AOJ
zCZ%tTevHuF{}ZEs5c&rnkO4?SAsZ70l?BWIVdNz6Lge;VK>cU+ZyEx#0xiBCZ`iPA
zDB6r;;iXg`Vah)`VtVU0pTA(^NkZaCKQ<M2q=01Ic>lGDB=@9pZX~A(1m3g>l$_py
z!2151k^F8z`RBd=xG@M*1_79kuq_f;mj(gRbNoJ5rdXIXCnBzZ0RnY~%U<WLIo$*r
zht1W*$C52ZB`X{~jgCLq|MG?OSo4o9av_mE{OA128Sh6R;0bjDf%c}8%K2TuZ;+E>
zfNzN5oGb^4tUwyAM`F;r(}(ll{Zh6YFRd{X*^(uZ9Mfbp>UTH=(NbYP-bliwghz@l
z521<ZpelRIx6Ouho5}!6|17z($TE~*h>0)>lL=m2^WZm)*dHJS@5xKv!IdJ<v^?62
z3t~n42u91DWGv)=w?ZgE9@?QY4YiPlJ}R?ASs^Fd%>SCZqln@u#E^J7Ls>T~SXww`
z>EnA@8Ctb&hq<v2QVX#-5cu3fjbd#NYiuO!KEiia5kD(nB!Gan(FeU08UrOSbeWkf
zp6@4=+(n!^u<Si#dZY@qINy2A<II{>#HW;9`t;~wzF;`NlU@I%LvY;T+U2LG?^@!_
zQc1LOPju~y!yO~>j11y<B%Zo@eQCkd(x~bEWKX89`NXJJorEdoxM3g{`Of!-&OlXp
zz&0#-#^huCdEEw?YSxTgsfE){_`ahfcXV<UH7_ppa;k(gCd*~B^;MHhudJ0y=P)*5
zaJoMpW42oepJ6;s&Yq0i5vCCuNK_%dWH5prFf8@uqgjr39Gaf)9l@b7Gq{{*WFZ9v
z)OaUqI|K$RQ~9jwMQav4#(nF$6f4}UHbRRlOwONX<k+=%=DQ9aYX0-RWM`AE=B?d_
z_t5H_KcBSso~wsy+FBv1_LZ!s1lYx76jK7+eBONB32T_Yr|6De(|SK7(X`xZfHLoF
z>79$qME?Z@bjGzoK+k%-dJ5KTU|><3Ctx-rEDuo0Dd|X}L$?Bb<odb}2t>(Ot8c{U
z`#<R;J1@l$SXy&I5j8R=>KaA^Q*PIk1YX_02U9+(@|@lBWqEBi@SvIoy~u~RlS-e1
zaMVt9W;*tq&<%1NWSGj`m^+dtc&pquk^6aNC2H)-sfA}TyR4n{D2_JG%cM28>A_b?
zof!9U+FWRVM;kjlwrL$ZJmU`Eczd`%i>YI}rH}BrCx;&65}uKy(>k%etmV)tsb?~$
zTDS4`<h~f4_bPa3HYdhW=(j?JOpFFb?Jvdpy+!tNHC{#6u|yJJI#;n0MOT!J_OXsc
z7P}_!(2a}X)FluY^W>fQs8~f=!n(7mYrVeGb5b19^T7xCqr)7kf3cIxy~sKQ_K!=d
z>ih6EgJIEMF4dLR_#D-~FoGXdNynmn(yDq(;%wADTjpyzR1tkNg=vJaM*lX3{c<3u
zi-H&a%AX%zGuE3RqfQ@=+r}tjq-Sb1ewC}fx($Bi5E6JK_jrq5m&&R?r44237ZNO&
z80jxW3H0?&99X~1599<|zwW8tH#2`3TQIGB<xT1l^t0n%c3#^OmK0uL;lV&c(r<{M
z+L60Rm(Py|2+2#~3UiZ*SEEoi4_86pHpM5aGUUKV4?&?B+Ps6HUr?<xsf}cksQAk2
zONeY6Z=OtUN7I8cG*0bMS#;B3iqdORm4#K(fLb<!nATo2%b1P6XMZ=z-^O98N^pF=
z*U!vdkW($s_w>%~&f&`7+H4iHe0_#vU4nAcJ=lX6<6Yi-!jU@_6LcEW&A*^pp2pD<
z=Z9>bS7{W6>KER8+h8BTuJroS<@9Opun$l7Y7Gb|(6)0GZL2$I2Yp{PX=YT;na`i1
zk>bWH-jYmIo~u}zg<_3XB>r9q0>14}oqXwetELXiAs~Qnv(qf0EeTUQe$)P@hWJq7
zFONYWNiOUb2vnd!z%Km2ZL=qben=)kSF1tb-WmG16je=-Bozt*=T(7-%x@V~4)Q_y
z_Rud<Elj;*oyXG#R*PTKoL(gde^K|0O&a)He`%~)yT!r@xBW9`x8=RaUiR@hFg6}(
z__4pT+w?A$VQd^;<UW?<IB2PLB3J7d;m09(m|C$jHcg5ieL2&nD^1v~)UaCqq_|Zy
zfPd7qJ$6*4KDt!r>7UZc<`-O*@+sFyz2~p-zT(2!CEKkqYVCP6LMt|aKTe+Nn<bD~
z^HD+xZrh`BqGAwE13~orTU~wM3=dy=IR|%I<pJEzxPgD=rbV&N-Gp{Bb!tV>is~ac
z^4g`y-E=}I4|Qa*40~-p7QLpozORhKP94#jWnLS-&_PHW9PC+p&&a52g&WJd22&)x
z<89Q-d}I>>#99yL+Ujc-GRu1@D8JB>xU*aB9E|6Qx+$&Y{VV0^$P*22nQMx|))24d
z^jBv%Q{#!2ijh|o)0lI0P=A>R!&GbCBe^Vx)2>GuxMS+?QfP#+UN5CY4AL{$1nKEf
zoM3g)&el2`5t5U<zFb{N9$QLEZ|8fhX@@EBPx)rIa)#b99^=!b2+ZnKjN&5_jBLU7
z*oLTvJ^OGJ?O>AQ=y|BpqUQar=AH#>wwKUrhK{)3T*8>1|4O23ruz?caw)aS4>M(G
z=U})nTWnz!lT>+TM;TXDkfZ8ioTd8mu1D;EYJK6Jq&M@!F_NcUi5GY)Q^b-#P>VKM
z2N{lZo9Q;{LdcE9?xE*h^sB_ap-Vk4xdsKpe^kagN{P>#Bi!6&r^UI&C1D_N%rLX<
ztIE~o)W-Oy)Ks{mow!oj7vU1IkHIvUB(66(GsW2&0z6~E3%SsDlX_`tG4fufbeJip
zx_@}^UiN@<P{5}?ueoA^zxC6x$NJ$TXJ>v*Cx)-=|CQs5Adl}fWwCmTU9jvof^!~O
zqwRT_If*v&h}74@PiF98N<|q;*Xg?EY>Ah**ErWPi%T7VIeUEW6%8Ku@Y_DxdU3Bq
z-EEP$O1YGCa<mFN+p;OFhGSLW%^XIG&LG>}*nqjlB+oAD(=0iewQkycj9HzQ9P)-D
zL%#b-dF6vq*GkMpluSTXTlz$`?*om!N8pkcD)M@TVb+up%Gh5`TWKJmJhN!TytQn|
zH!Z+$K3`X?R_A~~-|Xzo;lwjs|Jeu!QZfEQe3s?m4@%mvN#+@{XvQ<EVr%<?knY_R
zgz=B-gPFfmkw~^T+il@%uUp|NW#TJ!EgpMbkR;qWbu6!}u&J0|I8F+D*m9ohHM{?U
z#S8L%$G#T{hOyD47~1=^VEEq`bjS=_*HTJ`;Fxm&?m~|vioLWu9Ute-)Cy(t0(<65
zdh`{S1NpA_T@9)7<ci(9X8KGQeS_%sC*!WaCREn`ac9}dR@C_V@i+aDD<AZoB4a_o
z>>Y|s#*^<TusTulDfu(gO9Pc$li~b#HZjlTC>)!5qI6RshWuvb_1V?fdT-*Gr8o!}
z7%Q9}bnTGDt7%lu182KpuqXt|gz4AY;&b6st8>nwNUE>zpaF*6iO^#&O%eV5*y%PY
zBteWBT0_^R>-f2m4IMdjy_cGoEx#maxWfv!w!xquAIx$D0rR(rY|r|dS9PUN@S?ZR
zAu7J!!hY<J@FFt>$jD&rD9o!@MZf2yBv>^MsS1uV9?o7})4xIO{zNYb)89V8oF=>6
zZ%T2ZT^~f6r_Ob@=-yngEje)6Cr%}$>Ga8(TvgS>(;315%5XUfk9$vPwTYkec2Dlg
z9Au@3Qi`hjx?I+Gb2RQVmzFaX2ws}mMYRPe66KLrHdD_&{l0^QHXNRu@m#E4XC(nH
z`|h=_d~<K|IuGGLnb>}_oRsEh$LseY&|&^7SSHFO?^D)OF+dLZ4OXMf@e5pkkv*Cz
zXfaz&oYI4=-}#=Izf0VWKATGSf2owadF#9OfPlJJyV>ju?9i&QI&^XM$5i!yZ!P~{
rpOv2*^ncR^{pZh+e=4Yd&dLv|!k^3cXQ2EHl%KivUzb~@!M^_jdD0iu

diff --git a/examples/screenshots/webgpu_materials.jpg b/examples/screenshots/webgpu_materials.jpg
index 35b06935974bb9787ab26daee5a8712bd8be7e9b..001fe203e0b94e476d12e5c1c76aa979a0fbc77f 100644
GIT binary patch
delta 21
dcmZ2ya#d=>CFV^++8mn`ghCiu{xkf)2>@O+2%`W1

literal 6702
zcmeHLX;c$iwyqE$M&JU<prXxHL2zJ@84+<rhK^i7KnQ_W!5AaZT$x7)RLljDDF}!t
zqYNP;B#{sZAR>dvWtKn~0>~T^1QLd%9_U;B=e-~Hd#n3TyK9|&s_L9t^?m2;v+Fwr
zeg=mD87p&3a{vMX0LaDxfc-!eAOhL6O)s)(qR>q*CMpUQ6&Dj1|5jk|E#fdZOk8}6
z#1{D0P1`t=l-Me{dGY-*5iuxKY%5G0_O1D!a$pT0y9HQ<Jb^-F0TEdUR2Bl(L$Wq7
zihYB&0ra~;M4+N#;;;<}5;7aGq0kMmVjD;|dWLVb0iv>E+x8wkCBEJEGHjo(oX(y2
zms|FqF0GZfYv-uwUitkle5-=uj-9(y52&dhJfx>@U}$7~?9AD7=I1Ret?Vy4I67T&
zcDd@|c@2FXbK{nue?VYRaLB!=`wt#QKf)y>CMDxjo~EW{W#{DP<-aN@d{b6VB2(U0
zQ0v~+zi;@^*woz7+11_C`?-%X!WtbL|1vQ-HO*c4y12B=<FBkj0O)r*8~%Ur*&q%P
z*&K7&CLf4M;6{PUii+($D!%QME$p)I_I)~cw#c21e_2`!->+-Ok-zeL`&I=Ny%AOJ
zCZ%tTevHuF{}ZEs5c&rnkO4?SAsZ70l?BWIVdNz6Lge;VK>cU+ZyEx#0xiBCZ`iPA
zDB6r;;iXg`Vah)`VtVU0pTA(^NkZaCKQ<M2q=01Ic>lGDB=@9pZX~A(1m3g>l$_py
z!2151k^F8z`RBd=xG@M*1_79kuq_f;mj(gRbNoJ5rdXIXCnBzZ0RnY~%U<WLIo$*r
zht1W*$C52ZB`X{~jgCLq|MG?OSo4o9av_mE{OA128Sh6R;0bjDf%c}8%K2TuZ;+E>
zfNzN5oGb^4tUwyAM`F;r(}(ll{Zh6YFRd{X*^(uZ9Mfbp>UTH=(NbYP-bliwghz@l
z521<ZpelRIx6Ouho5}!6|17z($TE~*h>0)>lL=m2^WZm)*dHJS@5xKv!IdJ<v^?62
z3t~n42u91DWGv)=w?ZgE9@?QY4YiPlJ}R?ASs^Fd%>SCZqln@u#E^J7Ls>T~SXww`
z>EnA@8Ctb&hq<v2QVX#-5cu3fjbd#NYiuO!KEiia5kD(nB!Gan(FeU08UrOSbeWkf
zp6@4=+(n!^u<Si#dZY@qINy2A<II{>#HW;9`t;~wzF;`NlU@I%LvY;T+U2LG?^@!_
zQc1LOPju~y!yO~>j11y<B%Zo@eQCkd(x~bEWKX89`NXJJorEdoxM3g{`Of!-&OlXp
zz&0#-#^huCdEEw?YSxTgsfE){_`ahfcXV<UH7_ppa;k(gCd*~B^;MHhudJ0y=P)*5
zaJoMpW42oepJ6;s&Yq0i5vCCuNK_%dWH5prFf8@uqgjr39Gaf)9l@b7Gq{{*WFZ9v
z)OaUqI|K$RQ~9jwMQav4#(nF$6f4}UHbRRlOwONX<k+=%=DQ9aYX0-RWM`AE=B?d_
z_t5H_KcBSso~wsy+FBv1_LZ!s1lYx76jK7+eBONB32T_Yr|6De(|SK7(X`xZfHLoF
z>79$qME?Z@bjGzoK+k%-dJ5KTU|><3Ctx-rEDuo0Dd|X}L$?Bb<odb}2t>(Ot8c{U
z`#<R;J1@l$SXy&I5j8R=>KaA^Q*PIk1YX_02U9+(@|@lBWqEBi@SvIoy~u~RlS-e1
zaMVt9W;*tq&<%1NWSGj`m^+dtc&pquk^6aNC2H)-sfA}TyR4n{D2_JG%cM28>A_b?
zof!9U+FWRVM;kjlwrL$ZJmU`Eczd`%i>YI}rH}BrCx;&65}uKy(>k%etmV)tsb?~$
zTDS4`<h~f4_bPa3HYdhW=(j?JOpFFb?Jvdpy+!tNHC{#6u|yJJI#;n0MOT!J_OXsc
z7P}_!(2a}X)FluY^W>fQs8~f=!n(7mYrVeGb5b19^T7xCqr)7kf3cIxy~sKQ_K!=d
z>ih6EgJIEMF4dLR_#D-~FoGXdNynmn(yDq(;%wADTjpyzR1tkNg=vJaM*lX3{c<3u
zi-H&a%AX%zGuE3RqfQ@=+r}tjq-Sb1ewC}fx($Bi5E6JK_jrq5m&&R?r44237ZNO&
z80jxW3H0?&99X~1599<|zwW8tH#2`3TQIGB<xT1l^t0n%c3#^OmK0uL;lV&c(r<{M
z+L60Rm(Py|2+2#~3UiZ*SEEoi4_86pHpM5aGUUKV4?&?B+Ps6HUr?<xsf}cksQAk2
zONeY6Z=OtUN7I8cG*0bMS#;B3iqdORm4#K(fLb<!nATo2%b1P6XMZ=z-^O98N^pF=
z*U!vdkW($s_w>%~&f&`7+H4iHe0_#vU4nAcJ=lX6<6Yi-!jU@_6LcEW&A*^pp2pD<
z=Z9>bS7{W6>KER8+h8BTuJroS<@9Opun$l7Y7Gb|(6)0GZL2$I2Yp{PX=YT;na`i1
zk>bWH-jYmIo~u}zg<_3XB>r9q0>14}oqXwetELXiAs~Qnv(qf0EeTUQe$)P@hWJq7
zFONYWNiOUb2vnd!z%Km2ZL=qben=)kSF1tb-WmG16je=-Bozt*=T(7-%x@V~4)Q_y
z_Rud<Elj;*oyXG#R*PTKoL(gde^K|0O&a)He`%~)yT!r@xBW9`x8=RaUiR@hFg6}(
z__4pT+w?A$VQd^;<UW?<IB2PLB3J7d;m09(m|C$jHcg5ieL2&nD^1v~)UaCqq_|Zy
zfPd7qJ$6*4KDt!r>7UZc<`-O*@+sFyz2~p-zT(2!CEKkqYVCP6LMt|aKTe+Nn<bD~
z^HD+xZrh`BqGAwE13~orTU~wM3=dy=IR|%I<pJEzxPgD=rbV&N-Gp{Bb!tV>is~ac
z^4g`y-E=}I4|Qa*40~-p7QLpozORhKP94#jWnLS-&_PHW9PC+p&&a52g&WJd22&)x
z<89Q-d}I>>#99yL+Ujc-GRu1@D8JB>xU*aB9E|6Qx+$&Y{VV0^$P*22nQMx|))24d
z^jBv%Q{#!2ijh|o)0lI0P=A>R!&GbCBe^Vx)2>GuxMS+?QfP#+UN5CY4AL{$1nKEf
zoM3g)&el2`5t5U<zFb{N9$QLEZ|8fhX@@EBPx)rIa)#b99^=!b2+ZnKjN&5_jBLU7
z*oLTvJ^OGJ?O>AQ=y|BpqUQar=AH#>wwKUrhK{)3T*8>1|4O23ruz?caw)aS4>M(G
z=U})nTWnz!lT>+TM;TXDkfZ8ioTd8mu1D;EYJK6Jq&M@!F_NcUi5GY)Q^b-#P>VKM
z2N{lZo9Q;{LdcE9?xE*h^sB_ap-Vk4xdsKpe^kagN{P>#Bi!6&r^UI&C1D_N%rLX<
ztIE~o)W-Oy)Ks{mow!oj7vU1IkHIvUB(66(GsW2&0z6~E3%SsDlX_`tG4fufbeJip
zx_@}^UiN@<P{5}?ueoA^zxC6x$NJ$TXJ>v*Cx)-=|CQs5Adl}fWwCmTU9jvof^!~O
zqwRT_If*v&h}74@PiF98N<|q;*Xg?EY>Ah**ErWPi%T7VIeUEW6%8Ku@Y_DxdU3Bq
z-EEP$O1YGCa<mFN+p;OFhGSLW%^XIG&LG>}*nqjlB+oAD(=0iewQkycj9HzQ9P)-D
zL%#b-dF6vq*GkMpluSTXTlz$`?*om!N8pkcD)M@TVb+up%Gh5`TWKJmJhN!TytQn|
zH!Z+$K3`X?R_A~~-|Xzo;lwjs|Jeu!QZfEQe3s?m4@%mvN#+@{XvQ<EVr%<?knY_R
zgz=B-gPFfmkw~^T+il@%uUp|NW#TJ!EgpMbkR;qWbu6!}u&J0|I8F+D*m9ohHM{?U
z#S8L%$G#T{hOyD47~1=^VEEq`bjS=_*HTJ`;Fxm&?m~|vioLWu9Ute-)Cy(t0(<65
zdh`{S1NpA_T@9)7<ci(9X8KGQeS_%sC*!WaCREn`ac9}dR@C_V@i+aDD<AZoB4a_o
z>>Y|s#*^<TusTulDfu(gO9Pc$li~b#HZjlTC>)!5qI6RshWuvb_1V?fdT-*Gr8o!}
z7%Q9}bnTGDt7%lu182KpuqXt|gz4AY;&b6st8>nwNUE>zpaF*6iO^#&O%eV5*y%PY
zBteWBT0_^R>-f2m4IMdjy_cGoEx#maxWfv!w!xquAIx$D0rR(rY|r|dS9PUN@S?ZR
zAu7J!!hY<J@FFt>$jD&rD9o!@MZf2yBv>^MsS1uV9?o7})4xIO{zNYb)89V8oF=>6
zZ%T2ZT^~f6r_Ob@=-yngEje)6Cr%}$>Ga8(TvgS>(;315%5XUfk9$vPwTYkec2Dlg
z9Au@3Qi`hjx?I+Gb2RQVmzFaX2ws}mMYRPe66KLrHdD_&{l0^QHXNRu@m#E4XC(nH
z`|h=_d~<K|IuGGLnb>}_oRsEh$LseY&|&^7SSHFO?^D)OF+dLZ4OXMf@e5pkkv*Cz
zXfaz&oYI4=-}#=Izf0VWKATGSf2owadF#9OfPlJJyV>ju?9i&QI&^XM$5i!yZ!P~{
rpOv2*^ncR^{pZh+e=4Yd&dLv|!k^3cXQ2EHl%KivUzb~@!M^_jdD0iu

diff --git a/examples/screenshots/webgpu_nodes_playground.jpg b/examples/screenshots/webgpu_nodes_playground.jpg
index 35b06935974bb9787ab26daee5a8712bd8be7e9b..001fe203e0b94e476d12e5c1c76aa979a0fbc77f 100644
GIT binary patch
delta 21
dcmZ2ya#d=>CFV^++8mn`ghCiu{xkf)2>@O+2%`W1

literal 6702
zcmeHLX;c$iwyqE$M&JU<prXxHL2zJ@84+<rhK^i7KnQ_W!5AaZT$x7)RLljDDF}!t
zqYNP;B#{sZAR>dvWtKn~0>~T^1QLd%9_U;B=e-~Hd#n3TyK9|&s_L9t^?m2;v+Fwr
zeg=mD87p&3a{vMX0LaDxfc-!eAOhL6O)s)(qR>q*CMpUQ6&Dj1|5jk|E#fdZOk8}6
z#1{D0P1`t=l-Me{dGY-*5iuxKY%5G0_O1D!a$pT0y9HQ<Jb^-F0TEdUR2Bl(L$Wq7
zihYB&0ra~;M4+N#;;;<}5;7aGq0kMmVjD;|dWLVb0iv>E+x8wkCBEJEGHjo(oX(y2
zms|FqF0GZfYv-uwUitkle5-=uj-9(y52&dhJfx>@U}$7~?9AD7=I1Ret?Vy4I67T&
zcDd@|c@2FXbK{nue?VYRaLB!=`wt#QKf)y>CMDxjo~EW{W#{DP<-aN@d{b6VB2(U0
zQ0v~+zi;@^*woz7+11_C`?-%X!WtbL|1vQ-HO*c4y12B=<FBkj0O)r*8~%Ur*&q%P
z*&K7&CLf4M;6{PUii+($D!%QME$p)I_I)~cw#c21e_2`!->+-Ok-zeL`&I=Ny%AOJ
zCZ%tTevHuF{}ZEs5c&rnkO4?SAsZ70l?BWIVdNz6Lge;VK>cU+ZyEx#0xiBCZ`iPA
zDB6r;;iXg`Vah)`VtVU0pTA(^NkZaCKQ<M2q=01Ic>lGDB=@9pZX~A(1m3g>l$_py
z!2151k^F8z`RBd=xG@M*1_79kuq_f;mj(gRbNoJ5rdXIXCnBzZ0RnY~%U<WLIo$*r
zht1W*$C52ZB`X{~jgCLq|MG?OSo4o9av_mE{OA128Sh6R;0bjDf%c}8%K2TuZ;+E>
zfNzN5oGb^4tUwyAM`F;r(}(ll{Zh6YFRd{X*^(uZ9Mfbp>UTH=(NbYP-bliwghz@l
z521<ZpelRIx6Ouho5}!6|17z($TE~*h>0)>lL=m2^WZm)*dHJS@5xKv!IdJ<v^?62
z3t~n42u91DWGv)=w?ZgE9@?QY4YiPlJ}R?ASs^Fd%>SCZqln@u#E^J7Ls>T~SXww`
z>EnA@8Ctb&hq<v2QVX#-5cu3fjbd#NYiuO!KEiia5kD(nB!Gan(FeU08UrOSbeWkf
zp6@4=+(n!^u<Si#dZY@qINy2A<II{>#HW;9`t;~wzF;`NlU@I%LvY;T+U2LG?^@!_
zQc1LOPju~y!yO~>j11y<B%Zo@eQCkd(x~bEWKX89`NXJJorEdoxM3g{`Of!-&OlXp
zz&0#-#^huCdEEw?YSxTgsfE){_`ahfcXV<UH7_ppa;k(gCd*~B^;MHhudJ0y=P)*5
zaJoMpW42oepJ6;s&Yq0i5vCCuNK_%dWH5prFf8@uqgjr39Gaf)9l@b7Gq{{*WFZ9v
z)OaUqI|K$RQ~9jwMQav4#(nF$6f4}UHbRRlOwONX<k+=%=DQ9aYX0-RWM`AE=B?d_
z_t5H_KcBSso~wsy+FBv1_LZ!s1lYx76jK7+eBONB32T_Yr|6De(|SK7(X`xZfHLoF
z>79$qME?Z@bjGzoK+k%-dJ5KTU|><3Ctx-rEDuo0Dd|X}L$?Bb<odb}2t>(Ot8c{U
z`#<R;J1@l$SXy&I5j8R=>KaA^Q*PIk1YX_02U9+(@|@lBWqEBi@SvIoy~u~RlS-e1
zaMVt9W;*tq&<%1NWSGj`m^+dtc&pquk^6aNC2H)-sfA}TyR4n{D2_JG%cM28>A_b?
zof!9U+FWRVM;kjlwrL$ZJmU`Eczd`%i>YI}rH}BrCx;&65}uKy(>k%etmV)tsb?~$
zTDS4`<h~f4_bPa3HYdhW=(j?JOpFFb?Jvdpy+!tNHC{#6u|yJJI#;n0MOT!J_OXsc
z7P}_!(2a}X)FluY^W>fQs8~f=!n(7mYrVeGb5b19^T7xCqr)7kf3cIxy~sKQ_K!=d
z>ih6EgJIEMF4dLR_#D-~FoGXdNynmn(yDq(;%wADTjpyzR1tkNg=vJaM*lX3{c<3u
zi-H&a%AX%zGuE3RqfQ@=+r}tjq-Sb1ewC}fx($Bi5E6JK_jrq5m&&R?r44237ZNO&
z80jxW3H0?&99X~1599<|zwW8tH#2`3TQIGB<xT1l^t0n%c3#^OmK0uL;lV&c(r<{M
z+L60Rm(Py|2+2#~3UiZ*SEEoi4_86pHpM5aGUUKV4?&?B+Ps6HUr?<xsf}cksQAk2
zONeY6Z=OtUN7I8cG*0bMS#;B3iqdORm4#K(fLb<!nATo2%b1P6XMZ=z-^O98N^pF=
z*U!vdkW($s_w>%~&f&`7+H4iHe0_#vU4nAcJ=lX6<6Yi-!jU@_6LcEW&A*^pp2pD<
z=Z9>bS7{W6>KER8+h8BTuJroS<@9Opun$l7Y7Gb|(6)0GZL2$I2Yp{PX=YT;na`i1
zk>bWH-jYmIo~u}zg<_3XB>r9q0>14}oqXwetELXiAs~Qnv(qf0EeTUQe$)P@hWJq7
zFONYWNiOUb2vnd!z%Km2ZL=qben=)kSF1tb-WmG16je=-Bozt*=T(7-%x@V~4)Q_y
z_Rud<Elj;*oyXG#R*PTKoL(gde^K|0O&a)He`%~)yT!r@xBW9`x8=RaUiR@hFg6}(
z__4pT+w?A$VQd^;<UW?<IB2PLB3J7d;m09(m|C$jHcg5ieL2&nD^1v~)UaCqq_|Zy
zfPd7qJ$6*4KDt!r>7UZc<`-O*@+sFyz2~p-zT(2!CEKkqYVCP6LMt|aKTe+Nn<bD~
z^HD+xZrh`BqGAwE13~orTU~wM3=dy=IR|%I<pJEzxPgD=rbV&N-Gp{Bb!tV>is~ac
z^4g`y-E=}I4|Qa*40~-p7QLpozORhKP94#jWnLS-&_PHW9PC+p&&a52g&WJd22&)x
z<89Q-d}I>>#99yL+Ujc-GRu1@D8JB>xU*aB9E|6Qx+$&Y{VV0^$P*22nQMx|))24d
z^jBv%Q{#!2ijh|o)0lI0P=A>R!&GbCBe^Vx)2>GuxMS+?QfP#+UN5CY4AL{$1nKEf
zoM3g)&el2`5t5U<zFb{N9$QLEZ|8fhX@@EBPx)rIa)#b99^=!b2+ZnKjN&5_jBLU7
z*oLTvJ^OGJ?O>AQ=y|BpqUQar=AH#>wwKUrhK{)3T*8>1|4O23ruz?caw)aS4>M(G
z=U})nTWnz!lT>+TM;TXDkfZ8ioTd8mu1D;EYJK6Jq&M@!F_NcUi5GY)Q^b-#P>VKM
z2N{lZo9Q;{LdcE9?xE*h^sB_ap-Vk4xdsKpe^kagN{P>#Bi!6&r^UI&C1D_N%rLX<
ztIE~o)W-Oy)Ks{mow!oj7vU1IkHIvUB(66(GsW2&0z6~E3%SsDlX_`tG4fufbeJip
zx_@}^UiN@<P{5}?ueoA^zxC6x$NJ$TXJ>v*Cx)-=|CQs5Adl}fWwCmTU9jvof^!~O
zqwRT_If*v&h}74@PiF98N<|q;*Xg?EY>Ah**ErWPi%T7VIeUEW6%8Ku@Y_DxdU3Bq
z-EEP$O1YGCa<mFN+p;OFhGSLW%^XIG&LG>}*nqjlB+oAD(=0iewQkycj9HzQ9P)-D
zL%#b-dF6vq*GkMpluSTXTlz$`?*om!N8pkcD)M@TVb+up%Gh5`TWKJmJhN!TytQn|
zH!Z+$K3`X?R_A~~-|Xzo;lwjs|Jeu!QZfEQe3s?m4@%mvN#+@{XvQ<EVr%<?knY_R
zgz=B-gPFfmkw~^T+il@%uUp|NW#TJ!EgpMbkR;qWbu6!}u&J0|I8F+D*m9ohHM{?U
z#S8L%$G#T{hOyD47~1=^VEEq`bjS=_*HTJ`;Fxm&?m~|vioLWu9Ute-)Cy(t0(<65
zdh`{S1NpA_T@9)7<ci(9X8KGQeS_%sC*!WaCREn`ac9}dR@C_V@i+aDD<AZoB4a_o
z>>Y|s#*^<TusTulDfu(gO9Pc$li~b#HZjlTC>)!5qI6RshWuvb_1V?fdT-*Gr8o!}
z7%Q9}bnTGDt7%lu182KpuqXt|gz4AY;&b6st8>nwNUE>zpaF*6iO^#&O%eV5*y%PY
zBteWBT0_^R>-f2m4IMdjy_cGoEx#maxWfv!w!xquAIx$D0rR(rY|r|dS9PUN@S?ZR
zAu7J!!hY<J@FFt>$jD&rD9o!@MZf2yBv>^MsS1uV9?o7})4xIO{zNYb)89V8oF=>6
zZ%T2ZT^~f6r_Ob@=-yngEje)6Cr%}$>Ga8(TvgS>(;315%5XUfk9$vPwTYkec2Dlg
z9Au@3Qi`hjx?I+Gb2RQVmzFaX2ws}mMYRPe66KLrHdD_&{l0^QHXNRu@m#E4XC(nH
z`|h=_d~<K|IuGGLnb>}_oRsEh$LseY&|&^7SSHFO?^D)OF+dLZ4OXMf@e5pkkv*Cz
zXfaz&oYI4=-}#=Izf0VWKATGSf2owadF#9OfPlJJyV>ju?9i&QI&^XM$5i!yZ!P~{
rpOv2*^ncR^{pZh+e=4Yd&dLv|!k^3cXQ2EHl%KivUzb~@!M^_jdD0iu

diff --git a/examples/screenshots/webgpu_rtt.jpg b/examples/screenshots/webgpu_rtt.jpg
index 35b06935974bb9787ab26daee5a8712bd8be7e9b..001fe203e0b94e476d12e5c1c76aa979a0fbc77f 100644
GIT binary patch
delta 21
dcmZ2ya#d=>CFV^++8mn`ghCiu{xkf)2>@O+2%`W1

literal 6702
zcmeHLX;c$iwyqE$M&JU<prXxHL2zJ@84+<rhK^i7KnQ_W!5AaZT$x7)RLljDDF}!t
zqYNP;B#{sZAR>dvWtKn~0>~T^1QLd%9_U;B=e-~Hd#n3TyK9|&s_L9t^?m2;v+Fwr
zeg=mD87p&3a{vMX0LaDxfc-!eAOhL6O)s)(qR>q*CMpUQ6&Dj1|5jk|E#fdZOk8}6
z#1{D0P1`t=l-Me{dGY-*5iuxKY%5G0_O1D!a$pT0y9HQ<Jb^-F0TEdUR2Bl(L$Wq7
zihYB&0ra~;M4+N#;;;<}5;7aGq0kMmVjD;|dWLVb0iv>E+x8wkCBEJEGHjo(oX(y2
zms|FqF0GZfYv-uwUitkle5-=uj-9(y52&dhJfx>@U}$7~?9AD7=I1Ret?Vy4I67T&
zcDd@|c@2FXbK{nue?VYRaLB!=`wt#QKf)y>CMDxjo~EW{W#{DP<-aN@d{b6VB2(U0
zQ0v~+zi;@^*woz7+11_C`?-%X!WtbL|1vQ-HO*c4y12B=<FBkj0O)r*8~%Ur*&q%P
z*&K7&CLf4M;6{PUii+($D!%QME$p)I_I)~cw#c21e_2`!->+-Ok-zeL`&I=Ny%AOJ
zCZ%tTevHuF{}ZEs5c&rnkO4?SAsZ70l?BWIVdNz6Lge;VK>cU+ZyEx#0xiBCZ`iPA
zDB6r;;iXg`Vah)`VtVU0pTA(^NkZaCKQ<M2q=01Ic>lGDB=@9pZX~A(1m3g>l$_py
z!2151k^F8z`RBd=xG@M*1_79kuq_f;mj(gRbNoJ5rdXIXCnBzZ0RnY~%U<WLIo$*r
zht1W*$C52ZB`X{~jgCLq|MG?OSo4o9av_mE{OA128Sh6R;0bjDf%c}8%K2TuZ;+E>
zfNzN5oGb^4tUwyAM`F;r(}(ll{Zh6YFRd{X*^(uZ9Mfbp>UTH=(NbYP-bliwghz@l
z521<ZpelRIx6Ouho5}!6|17z($TE~*h>0)>lL=m2^WZm)*dHJS@5xKv!IdJ<v^?62
z3t~n42u91DWGv)=w?ZgE9@?QY4YiPlJ}R?ASs^Fd%>SCZqln@u#E^J7Ls>T~SXww`
z>EnA@8Ctb&hq<v2QVX#-5cu3fjbd#NYiuO!KEiia5kD(nB!Gan(FeU08UrOSbeWkf
zp6@4=+(n!^u<Si#dZY@qINy2A<II{>#HW;9`t;~wzF;`NlU@I%LvY;T+U2LG?^@!_
zQc1LOPju~y!yO~>j11y<B%Zo@eQCkd(x~bEWKX89`NXJJorEdoxM3g{`Of!-&OlXp
zz&0#-#^huCdEEw?YSxTgsfE){_`ahfcXV<UH7_ppa;k(gCd*~B^;MHhudJ0y=P)*5
zaJoMpW42oepJ6;s&Yq0i5vCCuNK_%dWH5prFf8@uqgjr39Gaf)9l@b7Gq{{*WFZ9v
z)OaUqI|K$RQ~9jwMQav4#(nF$6f4}UHbRRlOwONX<k+=%=DQ9aYX0-RWM`AE=B?d_
z_t5H_KcBSso~wsy+FBv1_LZ!s1lYx76jK7+eBONB32T_Yr|6De(|SK7(X`xZfHLoF
z>79$qME?Z@bjGzoK+k%-dJ5KTU|><3Ctx-rEDuo0Dd|X}L$?Bb<odb}2t>(Ot8c{U
z`#<R;J1@l$SXy&I5j8R=>KaA^Q*PIk1YX_02U9+(@|@lBWqEBi@SvIoy~u~RlS-e1
zaMVt9W;*tq&<%1NWSGj`m^+dtc&pquk^6aNC2H)-sfA}TyR4n{D2_JG%cM28>A_b?
zof!9U+FWRVM;kjlwrL$ZJmU`Eczd`%i>YI}rH}BrCx;&65}uKy(>k%etmV)tsb?~$
zTDS4`<h~f4_bPa3HYdhW=(j?JOpFFb?Jvdpy+!tNHC{#6u|yJJI#;n0MOT!J_OXsc
z7P}_!(2a}X)FluY^W>fQs8~f=!n(7mYrVeGb5b19^T7xCqr)7kf3cIxy~sKQ_K!=d
z>ih6EgJIEMF4dLR_#D-~FoGXdNynmn(yDq(;%wADTjpyzR1tkNg=vJaM*lX3{c<3u
zi-H&a%AX%zGuE3RqfQ@=+r}tjq-Sb1ewC}fx($Bi5E6JK_jrq5m&&R?r44237ZNO&
z80jxW3H0?&99X~1599<|zwW8tH#2`3TQIGB<xT1l^t0n%c3#^OmK0uL;lV&c(r<{M
z+L60Rm(Py|2+2#~3UiZ*SEEoi4_86pHpM5aGUUKV4?&?B+Ps6HUr?<xsf}cksQAk2
zONeY6Z=OtUN7I8cG*0bMS#;B3iqdORm4#K(fLb<!nATo2%b1P6XMZ=z-^O98N^pF=
z*U!vdkW($s_w>%~&f&`7+H4iHe0_#vU4nAcJ=lX6<6Yi-!jU@_6LcEW&A*^pp2pD<
z=Z9>bS7{W6>KER8+h8BTuJroS<@9Opun$l7Y7Gb|(6)0GZL2$I2Yp{PX=YT;na`i1
zk>bWH-jYmIo~u}zg<_3XB>r9q0>14}oqXwetELXiAs~Qnv(qf0EeTUQe$)P@hWJq7
zFONYWNiOUb2vnd!z%Km2ZL=qben=)kSF1tb-WmG16je=-Bozt*=T(7-%x@V~4)Q_y
z_Rud<Elj;*oyXG#R*PTKoL(gde^K|0O&a)He`%~)yT!r@xBW9`x8=RaUiR@hFg6}(
z__4pT+w?A$VQd^;<UW?<IB2PLB3J7d;m09(m|C$jHcg5ieL2&nD^1v~)UaCqq_|Zy
zfPd7qJ$6*4KDt!r>7UZc<`-O*@+sFyz2~p-zT(2!CEKkqYVCP6LMt|aKTe+Nn<bD~
z^HD+xZrh`BqGAwE13~orTU~wM3=dy=IR|%I<pJEzxPgD=rbV&N-Gp{Bb!tV>is~ac
z^4g`y-E=}I4|Qa*40~-p7QLpozORhKP94#jWnLS-&_PHW9PC+p&&a52g&WJd22&)x
z<89Q-d}I>>#99yL+Ujc-GRu1@D8JB>xU*aB9E|6Qx+$&Y{VV0^$P*22nQMx|))24d
z^jBv%Q{#!2ijh|o)0lI0P=A>R!&GbCBe^Vx)2>GuxMS+?QfP#+UN5CY4AL{$1nKEf
zoM3g)&el2`5t5U<zFb{N9$QLEZ|8fhX@@EBPx)rIa)#b99^=!b2+ZnKjN&5_jBLU7
z*oLTvJ^OGJ?O>AQ=y|BpqUQar=AH#>wwKUrhK{)3T*8>1|4O23ruz?caw)aS4>M(G
z=U})nTWnz!lT>+TM;TXDkfZ8ioTd8mu1D;EYJK6Jq&M@!F_NcUi5GY)Q^b-#P>VKM
z2N{lZo9Q;{LdcE9?xE*h^sB_ap-Vk4xdsKpe^kagN{P>#Bi!6&r^UI&C1D_N%rLX<
ztIE~o)W-Oy)Ks{mow!oj7vU1IkHIvUB(66(GsW2&0z6~E3%SsDlX_`tG4fufbeJip
zx_@}^UiN@<P{5}?ueoA^zxC6x$NJ$TXJ>v*Cx)-=|CQs5Adl}fWwCmTU9jvof^!~O
zqwRT_If*v&h}74@PiF98N<|q;*Xg?EY>Ah**ErWPi%T7VIeUEW6%8Ku@Y_DxdU3Bq
z-EEP$O1YGCa<mFN+p;OFhGSLW%^XIG&LG>}*nqjlB+oAD(=0iewQkycj9HzQ9P)-D
zL%#b-dF6vq*GkMpluSTXTlz$`?*om!N8pkcD)M@TVb+up%Gh5`TWKJmJhN!TytQn|
zH!Z+$K3`X?R_A~~-|Xzo;lwjs|Jeu!QZfEQe3s?m4@%mvN#+@{XvQ<EVr%<?knY_R
zgz=B-gPFfmkw~^T+il@%uUp|NW#TJ!EgpMbkR;qWbu6!}u&J0|I8F+D*m9ohHM{?U
z#S8L%$G#T{hOyD47~1=^VEEq`bjS=_*HTJ`;Fxm&?m~|vioLWu9Ute-)Cy(t0(<65
zdh`{S1NpA_T@9)7<ci(9X8KGQeS_%sC*!WaCREn`ac9}dR@C_V@i+aDD<AZoB4a_o
z>>Y|s#*^<TusTulDfu(gO9Pc$li~b#HZjlTC>)!5qI6RshWuvb_1V?fdT-*Gr8o!}
z7%Q9}bnTGDt7%lu182KpuqXt|gz4AY;&b6st8>nwNUE>zpaF*6iO^#&O%eV5*y%PY
zBteWBT0_^R>-f2m4IMdjy_cGoEx#maxWfv!w!xquAIx$D0rR(rY|r|dS9PUN@S?ZR
zAu7J!!hY<J@FFt>$jD&rD9o!@MZf2yBv>^MsS1uV9?o7})4xIO{zNYb)89V8oF=>6
zZ%T2ZT^~f6r_Ob@=-yngEje)6Cr%}$>Ga8(TvgS>(;315%5XUfk9$vPwTYkec2Dlg
z9Au@3Qi`hjx?I+Gb2RQVmzFaX2ws}mMYRPe66KLrHdD_&{l0^QHXNRu@m#E4XC(nH
z`|h=_d~<K|IuGGLnb>}_oRsEh$LseY&|&^7SSHFO?^D)OF+dLZ4OXMf@e5pkkv*Cz
zXfaz&oYI4=-}#=Izf0VWKATGSf2owadF#9OfPlJJyV>ju?9i&QI&^XM$5i!yZ!P~{
rpOv2*^ncR^{pZh+e=4Yd&dLv|!k^3cXQ2EHl%KivUzb~@!M^_jdD0iu

diff --git a/examples/screenshots/webgpu_sandbox.jpg b/examples/screenshots/webgpu_sandbox.jpg
index 35b06935974bb9787ab26daee5a8712bd8be7e9b..001fe203e0b94e476d12e5c1c76aa979a0fbc77f 100644
GIT binary patch
delta 21
dcmZ2ya#d=>CFV^++8mn`ghCiu{xkf)2>@O+2%`W1

literal 6702
zcmeHLX;c$iwyqE$M&JU<prXxHL2zJ@84+<rhK^i7KnQ_W!5AaZT$x7)RLljDDF}!t
zqYNP;B#{sZAR>dvWtKn~0>~T^1QLd%9_U;B=e-~Hd#n3TyK9|&s_L9t^?m2;v+Fwr
zeg=mD87p&3a{vMX0LaDxfc-!eAOhL6O)s)(qR>q*CMpUQ6&Dj1|5jk|E#fdZOk8}6
z#1{D0P1`t=l-Me{dGY-*5iuxKY%5G0_O1D!a$pT0y9HQ<Jb^-F0TEdUR2Bl(L$Wq7
zihYB&0ra~;M4+N#;;;<}5;7aGq0kMmVjD;|dWLVb0iv>E+x8wkCBEJEGHjo(oX(y2
zms|FqF0GZfYv-uwUitkle5-=uj-9(y52&dhJfx>@U}$7~?9AD7=I1Ret?Vy4I67T&
zcDd@|c@2FXbK{nue?VYRaLB!=`wt#QKf)y>CMDxjo~EW{W#{DP<-aN@d{b6VB2(U0
zQ0v~+zi;@^*woz7+11_C`?-%X!WtbL|1vQ-HO*c4y12B=<FBkj0O)r*8~%Ur*&q%P
z*&K7&CLf4M;6{PUii+($D!%QME$p)I_I)~cw#c21e_2`!->+-Ok-zeL`&I=Ny%AOJ
zCZ%tTevHuF{}ZEs5c&rnkO4?SAsZ70l?BWIVdNz6Lge;VK>cU+ZyEx#0xiBCZ`iPA
zDB6r;;iXg`Vah)`VtVU0pTA(^NkZaCKQ<M2q=01Ic>lGDB=@9pZX~A(1m3g>l$_py
z!2151k^F8z`RBd=xG@M*1_79kuq_f;mj(gRbNoJ5rdXIXCnBzZ0RnY~%U<WLIo$*r
zht1W*$C52ZB`X{~jgCLq|MG?OSo4o9av_mE{OA128Sh6R;0bjDf%c}8%K2TuZ;+E>
zfNzN5oGb^4tUwyAM`F;r(}(ll{Zh6YFRd{X*^(uZ9Mfbp>UTH=(NbYP-bliwghz@l
z521<ZpelRIx6Ouho5}!6|17z($TE~*h>0)>lL=m2^WZm)*dHJS@5xKv!IdJ<v^?62
z3t~n42u91DWGv)=w?ZgE9@?QY4YiPlJ}R?ASs^Fd%>SCZqln@u#E^J7Ls>T~SXww`
z>EnA@8Ctb&hq<v2QVX#-5cu3fjbd#NYiuO!KEiia5kD(nB!Gan(FeU08UrOSbeWkf
zp6@4=+(n!^u<Si#dZY@qINy2A<II{>#HW;9`t;~wzF;`NlU@I%LvY;T+U2LG?^@!_
zQc1LOPju~y!yO~>j11y<B%Zo@eQCkd(x~bEWKX89`NXJJorEdoxM3g{`Of!-&OlXp
zz&0#-#^huCdEEw?YSxTgsfE){_`ahfcXV<UH7_ppa;k(gCd*~B^;MHhudJ0y=P)*5
zaJoMpW42oepJ6;s&Yq0i5vCCuNK_%dWH5prFf8@uqgjr39Gaf)9l@b7Gq{{*WFZ9v
z)OaUqI|K$RQ~9jwMQav4#(nF$6f4}UHbRRlOwONX<k+=%=DQ9aYX0-RWM`AE=B?d_
z_t5H_KcBSso~wsy+FBv1_LZ!s1lYx76jK7+eBONB32T_Yr|6De(|SK7(X`xZfHLoF
z>79$qME?Z@bjGzoK+k%-dJ5KTU|><3Ctx-rEDuo0Dd|X}L$?Bb<odb}2t>(Ot8c{U
z`#<R;J1@l$SXy&I5j8R=>KaA^Q*PIk1YX_02U9+(@|@lBWqEBi@SvIoy~u~RlS-e1
zaMVt9W;*tq&<%1NWSGj`m^+dtc&pquk^6aNC2H)-sfA}TyR4n{D2_JG%cM28>A_b?
zof!9U+FWRVM;kjlwrL$ZJmU`Eczd`%i>YI}rH}BrCx;&65}uKy(>k%etmV)tsb?~$
zTDS4`<h~f4_bPa3HYdhW=(j?JOpFFb?Jvdpy+!tNHC{#6u|yJJI#;n0MOT!J_OXsc
z7P}_!(2a}X)FluY^W>fQs8~f=!n(7mYrVeGb5b19^T7xCqr)7kf3cIxy~sKQ_K!=d
z>ih6EgJIEMF4dLR_#D-~FoGXdNynmn(yDq(;%wADTjpyzR1tkNg=vJaM*lX3{c<3u
zi-H&a%AX%zGuE3RqfQ@=+r}tjq-Sb1ewC}fx($Bi5E6JK_jrq5m&&R?r44237ZNO&
z80jxW3H0?&99X~1599<|zwW8tH#2`3TQIGB<xT1l^t0n%c3#^OmK0uL;lV&c(r<{M
z+L60Rm(Py|2+2#~3UiZ*SEEoi4_86pHpM5aGUUKV4?&?B+Ps6HUr?<xsf}cksQAk2
zONeY6Z=OtUN7I8cG*0bMS#;B3iqdORm4#K(fLb<!nATo2%b1P6XMZ=z-^O98N^pF=
z*U!vdkW($s_w>%~&f&`7+H4iHe0_#vU4nAcJ=lX6<6Yi-!jU@_6LcEW&A*^pp2pD<
z=Z9>bS7{W6>KER8+h8BTuJroS<@9Opun$l7Y7Gb|(6)0GZL2$I2Yp{PX=YT;na`i1
zk>bWH-jYmIo~u}zg<_3XB>r9q0>14}oqXwetELXiAs~Qnv(qf0EeTUQe$)P@hWJq7
zFONYWNiOUb2vnd!z%Km2ZL=qben=)kSF1tb-WmG16je=-Bozt*=T(7-%x@V~4)Q_y
z_Rud<Elj;*oyXG#R*PTKoL(gde^K|0O&a)He`%~)yT!r@xBW9`x8=RaUiR@hFg6}(
z__4pT+w?A$VQd^;<UW?<IB2PLB3J7d;m09(m|C$jHcg5ieL2&nD^1v~)UaCqq_|Zy
zfPd7qJ$6*4KDt!r>7UZc<`-O*@+sFyz2~p-zT(2!CEKkqYVCP6LMt|aKTe+Nn<bD~
z^HD+xZrh`BqGAwE13~orTU~wM3=dy=IR|%I<pJEzxPgD=rbV&N-Gp{Bb!tV>is~ac
z^4g`y-E=}I4|Qa*40~-p7QLpozORhKP94#jWnLS-&_PHW9PC+p&&a52g&WJd22&)x
z<89Q-d}I>>#99yL+Ujc-GRu1@D8JB>xU*aB9E|6Qx+$&Y{VV0^$P*22nQMx|))24d
z^jBv%Q{#!2ijh|o)0lI0P=A>R!&GbCBe^Vx)2>GuxMS+?QfP#+UN5CY4AL{$1nKEf
zoM3g)&el2`5t5U<zFb{N9$QLEZ|8fhX@@EBPx)rIa)#b99^=!b2+ZnKjN&5_jBLU7
z*oLTvJ^OGJ?O>AQ=y|BpqUQar=AH#>wwKUrhK{)3T*8>1|4O23ruz?caw)aS4>M(G
z=U})nTWnz!lT>+TM;TXDkfZ8ioTd8mu1D;EYJK6Jq&M@!F_NcUi5GY)Q^b-#P>VKM
z2N{lZo9Q;{LdcE9?xE*h^sB_ap-Vk4xdsKpe^kagN{P>#Bi!6&r^UI&C1D_N%rLX<
ztIE~o)W-Oy)Ks{mow!oj7vU1IkHIvUB(66(GsW2&0z6~E3%SsDlX_`tG4fufbeJip
zx_@}^UiN@<P{5}?ueoA^zxC6x$NJ$TXJ>v*Cx)-=|CQs5Adl}fWwCmTU9jvof^!~O
zqwRT_If*v&h}74@PiF98N<|q;*Xg?EY>Ah**ErWPi%T7VIeUEW6%8Ku@Y_DxdU3Bq
z-EEP$O1YGCa<mFN+p;OFhGSLW%^XIG&LG>}*nqjlB+oAD(=0iewQkycj9HzQ9P)-D
zL%#b-dF6vq*GkMpluSTXTlz$`?*om!N8pkcD)M@TVb+up%Gh5`TWKJmJhN!TytQn|
zH!Z+$K3`X?R_A~~-|Xzo;lwjs|Jeu!QZfEQe3s?m4@%mvN#+@{XvQ<EVr%<?knY_R
zgz=B-gPFfmkw~^T+il@%uUp|NW#TJ!EgpMbkR;qWbu6!}u&J0|I8F+D*m9ohHM{?U
z#S8L%$G#T{hOyD47~1=^VEEq`bjS=_*HTJ`;Fxm&?m~|vioLWu9Ute-)Cy(t0(<65
zdh`{S1NpA_T@9)7<ci(9X8KGQeS_%sC*!WaCREn`ac9}dR@C_V@i+aDD<AZoB4a_o
z>>Y|s#*^<TusTulDfu(gO9Pc$li~b#HZjlTC>)!5qI6RshWuvb_1V?fdT-*Gr8o!}
z7%Q9}bnTGDt7%lu182KpuqXt|gz4AY;&b6st8>nwNUE>zpaF*6iO^#&O%eV5*y%PY
zBteWBT0_^R>-f2m4IMdjy_cGoEx#maxWfv!w!xquAIx$D0rR(rY|r|dS9PUN@S?ZR
zAu7J!!hY<J@FFt>$jD&rD9o!@MZf2yBv>^MsS1uV9?o7})4xIO{zNYb)89V8oF=>6
zZ%T2ZT^~f6r_Ob@=-yngEje)6Cr%}$>Ga8(TvgS>(;315%5XUfk9$vPwTYkec2Dlg
z9Au@3Qi`hjx?I+Gb2RQVmzFaX2ws}mMYRPe66KLrHdD_&{l0^QHXNRu@m#E4XC(nH
z`|h=_d~<K|IuGGLnb>}_oRsEh$LseY&|&^7SSHFO?^D)OF+dLZ4OXMf@e5pkkv*Cz
zXfaz&oYI4=-}#=Izf0VWKATGSf2owadF#9OfPlJJyV>ju?9i&QI&^XM$5i!yZ!P~{
rpOv2*^ncR^{pZh+e=4Yd&dLv|!k^3cXQ2EHl%KivUzb~@!M^_jdD0iu

diff --git a/examples/screenshots/webgpu_skinning.jpg b/examples/screenshots/webgpu_skinning.jpg
index 35b06935974bb9787ab26daee5a8712bd8be7e9b..001fe203e0b94e476d12e5c1c76aa979a0fbc77f 100644
GIT binary patch
delta 21
dcmZ2ya#d=>CFV^++8mn`ghCiu{xkf)2>@O+2%`W1

literal 6702
zcmeHLX;c$iwyqE$M&JU<prXxHL2zJ@84+<rhK^i7KnQ_W!5AaZT$x7)RLljDDF}!t
zqYNP;B#{sZAR>dvWtKn~0>~T^1QLd%9_U;B=e-~Hd#n3TyK9|&s_L9t^?m2;v+Fwr
zeg=mD87p&3a{vMX0LaDxfc-!eAOhL6O)s)(qR>q*CMpUQ6&Dj1|5jk|E#fdZOk8}6
z#1{D0P1`t=l-Me{dGY-*5iuxKY%5G0_O1D!a$pT0y9HQ<Jb^-F0TEdUR2Bl(L$Wq7
zihYB&0ra~;M4+N#;;;<}5;7aGq0kMmVjD;|dWLVb0iv>E+x8wkCBEJEGHjo(oX(y2
zms|FqF0GZfYv-uwUitkle5-=uj-9(y52&dhJfx>@U}$7~?9AD7=I1Ret?Vy4I67T&
zcDd@|c@2FXbK{nue?VYRaLB!=`wt#QKf)y>CMDxjo~EW{W#{DP<-aN@d{b6VB2(U0
zQ0v~+zi;@^*woz7+11_C`?-%X!WtbL|1vQ-HO*c4y12B=<FBkj0O)r*8~%Ur*&q%P
z*&K7&CLf4M;6{PUii+($D!%QME$p)I_I)~cw#c21e_2`!->+-Ok-zeL`&I=Ny%AOJ
zCZ%tTevHuF{}ZEs5c&rnkO4?SAsZ70l?BWIVdNz6Lge;VK>cU+ZyEx#0xiBCZ`iPA
zDB6r;;iXg`Vah)`VtVU0pTA(^NkZaCKQ<M2q=01Ic>lGDB=@9pZX~A(1m3g>l$_py
z!2151k^F8z`RBd=xG@M*1_79kuq_f;mj(gRbNoJ5rdXIXCnBzZ0RnY~%U<WLIo$*r
zht1W*$C52ZB`X{~jgCLq|MG?OSo4o9av_mE{OA128Sh6R;0bjDf%c}8%K2TuZ;+E>
zfNzN5oGb^4tUwyAM`F;r(}(ll{Zh6YFRd{X*^(uZ9Mfbp>UTH=(NbYP-bliwghz@l
z521<ZpelRIx6Ouho5}!6|17z($TE~*h>0)>lL=m2^WZm)*dHJS@5xKv!IdJ<v^?62
z3t~n42u91DWGv)=w?ZgE9@?QY4YiPlJ}R?ASs^Fd%>SCZqln@u#E^J7Ls>T~SXww`
z>EnA@8Ctb&hq<v2QVX#-5cu3fjbd#NYiuO!KEiia5kD(nB!Gan(FeU08UrOSbeWkf
zp6@4=+(n!^u<Si#dZY@qINy2A<II{>#HW;9`t;~wzF;`NlU@I%LvY;T+U2LG?^@!_
zQc1LOPju~y!yO~>j11y<B%Zo@eQCkd(x~bEWKX89`NXJJorEdoxM3g{`Of!-&OlXp
zz&0#-#^huCdEEw?YSxTgsfE){_`ahfcXV<UH7_ppa;k(gCd*~B^;MHhudJ0y=P)*5
zaJoMpW42oepJ6;s&Yq0i5vCCuNK_%dWH5prFf8@uqgjr39Gaf)9l@b7Gq{{*WFZ9v
z)OaUqI|K$RQ~9jwMQav4#(nF$6f4}UHbRRlOwONX<k+=%=DQ9aYX0-RWM`AE=B?d_
z_t5H_KcBSso~wsy+FBv1_LZ!s1lYx76jK7+eBONB32T_Yr|6De(|SK7(X`xZfHLoF
z>79$qME?Z@bjGzoK+k%-dJ5KTU|><3Ctx-rEDuo0Dd|X}L$?Bb<odb}2t>(Ot8c{U
z`#<R;J1@l$SXy&I5j8R=>KaA^Q*PIk1YX_02U9+(@|@lBWqEBi@SvIoy~u~RlS-e1
zaMVt9W;*tq&<%1NWSGj`m^+dtc&pquk^6aNC2H)-sfA}TyR4n{D2_JG%cM28>A_b?
zof!9U+FWRVM;kjlwrL$ZJmU`Eczd`%i>YI}rH}BrCx;&65}uKy(>k%etmV)tsb?~$
zTDS4`<h~f4_bPa3HYdhW=(j?JOpFFb?Jvdpy+!tNHC{#6u|yJJI#;n0MOT!J_OXsc
z7P}_!(2a}X)FluY^W>fQs8~f=!n(7mYrVeGb5b19^T7xCqr)7kf3cIxy~sKQ_K!=d
z>ih6EgJIEMF4dLR_#D-~FoGXdNynmn(yDq(;%wADTjpyzR1tkNg=vJaM*lX3{c<3u
zi-H&a%AX%zGuE3RqfQ@=+r}tjq-Sb1ewC}fx($Bi5E6JK_jrq5m&&R?r44237ZNO&
z80jxW3H0?&99X~1599<|zwW8tH#2`3TQIGB<xT1l^t0n%c3#^OmK0uL;lV&c(r<{M
z+L60Rm(Py|2+2#~3UiZ*SEEoi4_86pHpM5aGUUKV4?&?B+Ps6HUr?<xsf}cksQAk2
zONeY6Z=OtUN7I8cG*0bMS#;B3iqdORm4#K(fLb<!nATo2%b1P6XMZ=z-^O98N^pF=
z*U!vdkW($s_w>%~&f&`7+H4iHe0_#vU4nAcJ=lX6<6Yi-!jU@_6LcEW&A*^pp2pD<
z=Z9>bS7{W6>KER8+h8BTuJroS<@9Opun$l7Y7Gb|(6)0GZL2$I2Yp{PX=YT;na`i1
zk>bWH-jYmIo~u}zg<_3XB>r9q0>14}oqXwetELXiAs~Qnv(qf0EeTUQe$)P@hWJq7
zFONYWNiOUb2vnd!z%Km2ZL=qben=)kSF1tb-WmG16je=-Bozt*=T(7-%x@V~4)Q_y
z_Rud<Elj;*oyXG#R*PTKoL(gde^K|0O&a)He`%~)yT!r@xBW9`x8=RaUiR@hFg6}(
z__4pT+w?A$VQd^;<UW?<IB2PLB3J7d;m09(m|C$jHcg5ieL2&nD^1v~)UaCqq_|Zy
zfPd7qJ$6*4KDt!r>7UZc<`-O*@+sFyz2~p-zT(2!CEKkqYVCP6LMt|aKTe+Nn<bD~
z^HD+xZrh`BqGAwE13~orTU~wM3=dy=IR|%I<pJEzxPgD=rbV&N-Gp{Bb!tV>is~ac
z^4g`y-E=}I4|Qa*40~-p7QLpozORhKP94#jWnLS-&_PHW9PC+p&&a52g&WJd22&)x
z<89Q-d}I>>#99yL+Ujc-GRu1@D8JB>xU*aB9E|6Qx+$&Y{VV0^$P*22nQMx|))24d
z^jBv%Q{#!2ijh|o)0lI0P=A>R!&GbCBe^Vx)2>GuxMS+?QfP#+UN5CY4AL{$1nKEf
zoM3g)&el2`5t5U<zFb{N9$QLEZ|8fhX@@EBPx)rIa)#b99^=!b2+ZnKjN&5_jBLU7
z*oLTvJ^OGJ?O>AQ=y|BpqUQar=AH#>wwKUrhK{)3T*8>1|4O23ruz?caw)aS4>M(G
z=U})nTWnz!lT>+TM;TXDkfZ8ioTd8mu1D;EYJK6Jq&M@!F_NcUi5GY)Q^b-#P>VKM
z2N{lZo9Q;{LdcE9?xE*h^sB_ap-Vk4xdsKpe^kagN{P>#Bi!6&r^UI&C1D_N%rLX<
ztIE~o)W-Oy)Ks{mow!oj7vU1IkHIvUB(66(GsW2&0z6~E3%SsDlX_`tG4fufbeJip
zx_@}^UiN@<P{5}?ueoA^zxC6x$NJ$TXJ>v*Cx)-=|CQs5Adl}fWwCmTU9jvof^!~O
zqwRT_If*v&h}74@PiF98N<|q;*Xg?EY>Ah**ErWPi%T7VIeUEW6%8Ku@Y_DxdU3Bq
z-EEP$O1YGCa<mFN+p;OFhGSLW%^XIG&LG>}*nqjlB+oAD(=0iewQkycj9HzQ9P)-D
zL%#b-dF6vq*GkMpluSTXTlz$`?*om!N8pkcD)M@TVb+up%Gh5`TWKJmJhN!TytQn|
zH!Z+$K3`X?R_A~~-|Xzo;lwjs|Jeu!QZfEQe3s?m4@%mvN#+@{XvQ<EVr%<?knY_R
zgz=B-gPFfmkw~^T+il@%uUp|NW#TJ!EgpMbkR;qWbu6!}u&J0|I8F+D*m9ohHM{?U
z#S8L%$G#T{hOyD47~1=^VEEq`bjS=_*HTJ`;Fxm&?m~|vioLWu9Ute-)Cy(t0(<65
zdh`{S1NpA_T@9)7<ci(9X8KGQeS_%sC*!WaCREn`ac9}dR@C_V@i+aDD<AZoB4a_o
z>>Y|s#*^<TusTulDfu(gO9Pc$li~b#HZjlTC>)!5qI6RshWuvb_1V?fdT-*Gr8o!}
z7%Q9}bnTGDt7%lu182KpuqXt|gz4AY;&b6st8>nwNUE>zpaF*6iO^#&O%eV5*y%PY
zBteWBT0_^R>-f2m4IMdjy_cGoEx#maxWfv!w!xquAIx$D0rR(rY|r|dS9PUN@S?ZR
zAu7J!!hY<J@FFt>$jD&rD9o!@MZf2yBv>^MsS1uV9?o7})4xIO{zNYb)89V8oF=>6
zZ%T2ZT^~f6r_Ob@=-yngEje)6Cr%}$>Ga8(TvgS>(;315%5XUfk9$vPwTYkec2Dlg
z9Au@3Qi`hjx?I+Gb2RQVmzFaX2ws}mMYRPe66KLrHdD_&{l0^QHXNRu@m#E4XC(nH
z`|h=_d~<K|IuGGLnb>}_oRsEh$LseY&|&^7SSHFO?^D)OF+dLZ4OXMf@e5pkkv*Cz
zXfaz&oYI4=-}#=Izf0VWKATGSf2owadF#9OfPlJJyV>ju?9i&QI&^XM$5i!yZ!P~{
rpOv2*^ncR^{pZh+e=4Yd&dLv|!k^3cXQ2EHl%KivUzb~@!M^_jdD0iu

diff --git a/examples/screenshots/webgpu_skinning_instancing.jpg b/examples/screenshots/webgpu_skinning_instancing.jpg
index 35b06935974bb9787ab26daee5a8712bd8be7e9b..001fe203e0b94e476d12e5c1c76aa979a0fbc77f 100644
GIT binary patch
delta 21
dcmZ2ya#d=>CFV^++8mn`ghCiu{xkf)2>@O+2%`W1

literal 6702
zcmeHLX;c$iwyqE$M&JU<prXxHL2zJ@84+<rhK^i7KnQ_W!5AaZT$x7)RLljDDF}!t
zqYNP;B#{sZAR>dvWtKn~0>~T^1QLd%9_U;B=e-~Hd#n3TyK9|&s_L9t^?m2;v+Fwr
zeg=mD87p&3a{vMX0LaDxfc-!eAOhL6O)s)(qR>q*CMpUQ6&Dj1|5jk|E#fdZOk8}6
z#1{D0P1`t=l-Me{dGY-*5iuxKY%5G0_O1D!a$pT0y9HQ<Jb^-F0TEdUR2Bl(L$Wq7
zihYB&0ra~;M4+N#;;;<}5;7aGq0kMmVjD;|dWLVb0iv>E+x8wkCBEJEGHjo(oX(y2
zms|FqF0GZfYv-uwUitkle5-=uj-9(y52&dhJfx>@U}$7~?9AD7=I1Ret?Vy4I67T&
zcDd@|c@2FXbK{nue?VYRaLB!=`wt#QKf)y>CMDxjo~EW{W#{DP<-aN@d{b6VB2(U0
zQ0v~+zi;@^*woz7+11_C`?-%X!WtbL|1vQ-HO*c4y12B=<FBkj0O)r*8~%Ur*&q%P
z*&K7&CLf4M;6{PUii+($D!%QME$p)I_I)~cw#c21e_2`!->+-Ok-zeL`&I=Ny%AOJ
zCZ%tTevHuF{}ZEs5c&rnkO4?SAsZ70l?BWIVdNz6Lge;VK>cU+ZyEx#0xiBCZ`iPA
zDB6r;;iXg`Vah)`VtVU0pTA(^NkZaCKQ<M2q=01Ic>lGDB=@9pZX~A(1m3g>l$_py
z!2151k^F8z`RBd=xG@M*1_79kuq_f;mj(gRbNoJ5rdXIXCnBzZ0RnY~%U<WLIo$*r
zht1W*$C52ZB`X{~jgCLq|MG?OSo4o9av_mE{OA128Sh6R;0bjDf%c}8%K2TuZ;+E>
zfNzN5oGb^4tUwyAM`F;r(}(ll{Zh6YFRd{X*^(uZ9Mfbp>UTH=(NbYP-bliwghz@l
z521<ZpelRIx6Ouho5}!6|17z($TE~*h>0)>lL=m2^WZm)*dHJS@5xKv!IdJ<v^?62
z3t~n42u91DWGv)=w?ZgE9@?QY4YiPlJ}R?ASs^Fd%>SCZqln@u#E^J7Ls>T~SXww`
z>EnA@8Ctb&hq<v2QVX#-5cu3fjbd#NYiuO!KEiia5kD(nB!Gan(FeU08UrOSbeWkf
zp6@4=+(n!^u<Si#dZY@qINy2A<II{>#HW;9`t;~wzF;`NlU@I%LvY;T+U2LG?^@!_
zQc1LOPju~y!yO~>j11y<B%Zo@eQCkd(x~bEWKX89`NXJJorEdoxM3g{`Of!-&OlXp
zz&0#-#^huCdEEw?YSxTgsfE){_`ahfcXV<UH7_ppa;k(gCd*~B^;MHhudJ0y=P)*5
zaJoMpW42oepJ6;s&Yq0i5vCCuNK_%dWH5prFf8@uqgjr39Gaf)9l@b7Gq{{*WFZ9v
z)OaUqI|K$RQ~9jwMQav4#(nF$6f4}UHbRRlOwONX<k+=%=DQ9aYX0-RWM`AE=B?d_
z_t5H_KcBSso~wsy+FBv1_LZ!s1lYx76jK7+eBONB32T_Yr|6De(|SK7(X`xZfHLoF
z>79$qME?Z@bjGzoK+k%-dJ5KTU|><3Ctx-rEDuo0Dd|X}L$?Bb<odb}2t>(Ot8c{U
z`#<R;J1@l$SXy&I5j8R=>KaA^Q*PIk1YX_02U9+(@|@lBWqEBi@SvIoy~u~RlS-e1
zaMVt9W;*tq&<%1NWSGj`m^+dtc&pquk^6aNC2H)-sfA}TyR4n{D2_JG%cM28>A_b?
zof!9U+FWRVM;kjlwrL$ZJmU`Eczd`%i>YI}rH}BrCx;&65}uKy(>k%etmV)tsb?~$
zTDS4`<h~f4_bPa3HYdhW=(j?JOpFFb?Jvdpy+!tNHC{#6u|yJJI#;n0MOT!J_OXsc
z7P}_!(2a}X)FluY^W>fQs8~f=!n(7mYrVeGb5b19^T7xCqr)7kf3cIxy~sKQ_K!=d
z>ih6EgJIEMF4dLR_#D-~FoGXdNynmn(yDq(;%wADTjpyzR1tkNg=vJaM*lX3{c<3u
zi-H&a%AX%zGuE3RqfQ@=+r}tjq-Sb1ewC}fx($Bi5E6JK_jrq5m&&R?r44237ZNO&
z80jxW3H0?&99X~1599<|zwW8tH#2`3TQIGB<xT1l^t0n%c3#^OmK0uL;lV&c(r<{M
z+L60Rm(Py|2+2#~3UiZ*SEEoi4_86pHpM5aGUUKV4?&?B+Ps6HUr?<xsf}cksQAk2
zONeY6Z=OtUN7I8cG*0bMS#;B3iqdORm4#K(fLb<!nATo2%b1P6XMZ=z-^O98N^pF=
z*U!vdkW($s_w>%~&f&`7+H4iHe0_#vU4nAcJ=lX6<6Yi-!jU@_6LcEW&A*^pp2pD<
z=Z9>bS7{W6>KER8+h8BTuJroS<@9Opun$l7Y7Gb|(6)0GZL2$I2Yp{PX=YT;na`i1
zk>bWH-jYmIo~u}zg<_3XB>r9q0>14}oqXwetELXiAs~Qnv(qf0EeTUQe$)P@hWJq7
zFONYWNiOUb2vnd!z%Km2ZL=qben=)kSF1tb-WmG16je=-Bozt*=T(7-%x@V~4)Q_y
z_Rud<Elj;*oyXG#R*PTKoL(gde^K|0O&a)He`%~)yT!r@xBW9`x8=RaUiR@hFg6}(
z__4pT+w?A$VQd^;<UW?<IB2PLB3J7d;m09(m|C$jHcg5ieL2&nD^1v~)UaCqq_|Zy
zfPd7qJ$6*4KDt!r>7UZc<`-O*@+sFyz2~p-zT(2!CEKkqYVCP6LMt|aKTe+Nn<bD~
z^HD+xZrh`BqGAwE13~orTU~wM3=dy=IR|%I<pJEzxPgD=rbV&N-Gp{Bb!tV>is~ac
z^4g`y-E=}I4|Qa*40~-p7QLpozORhKP94#jWnLS-&_PHW9PC+p&&a52g&WJd22&)x
z<89Q-d}I>>#99yL+Ujc-GRu1@D8JB>xU*aB9E|6Qx+$&Y{VV0^$P*22nQMx|))24d
z^jBv%Q{#!2ijh|o)0lI0P=A>R!&GbCBe^Vx)2>GuxMS+?QfP#+UN5CY4AL{$1nKEf
zoM3g)&el2`5t5U<zFb{N9$QLEZ|8fhX@@EBPx)rIa)#b99^=!b2+ZnKjN&5_jBLU7
z*oLTvJ^OGJ?O>AQ=y|BpqUQar=AH#>wwKUrhK{)3T*8>1|4O23ruz?caw)aS4>M(G
z=U})nTWnz!lT>+TM;TXDkfZ8ioTd8mu1D;EYJK6Jq&M@!F_NcUi5GY)Q^b-#P>VKM
z2N{lZo9Q;{LdcE9?xE*h^sB_ap-Vk4xdsKpe^kagN{P>#Bi!6&r^UI&C1D_N%rLX<
ztIE~o)W-Oy)Ks{mow!oj7vU1IkHIvUB(66(GsW2&0z6~E3%SsDlX_`tG4fufbeJip
zx_@}^UiN@<P{5}?ueoA^zxC6x$NJ$TXJ>v*Cx)-=|CQs5Adl}fWwCmTU9jvof^!~O
zqwRT_If*v&h}74@PiF98N<|q;*Xg?EY>Ah**ErWPi%T7VIeUEW6%8Ku@Y_DxdU3Bq
z-EEP$O1YGCa<mFN+p;OFhGSLW%^XIG&LG>}*nqjlB+oAD(=0iewQkycj9HzQ9P)-D
zL%#b-dF6vq*GkMpluSTXTlz$`?*om!N8pkcD)M@TVb+up%Gh5`TWKJmJhN!TytQn|
zH!Z+$K3`X?R_A~~-|Xzo;lwjs|Jeu!QZfEQe3s?m4@%mvN#+@{XvQ<EVr%<?knY_R
zgz=B-gPFfmkw~^T+il@%uUp|NW#TJ!EgpMbkR;qWbu6!}u&J0|I8F+D*m9ohHM{?U
z#S8L%$G#T{hOyD47~1=^VEEq`bjS=_*HTJ`;Fxm&?m~|vioLWu9Ute-)Cy(t0(<65
zdh`{S1NpA_T@9)7<ci(9X8KGQeS_%sC*!WaCREn`ac9}dR@C_V@i+aDD<AZoB4a_o
z>>Y|s#*^<TusTulDfu(gO9Pc$li~b#HZjlTC>)!5qI6RshWuvb_1V?fdT-*Gr8o!}
z7%Q9}bnTGDt7%lu182KpuqXt|gz4AY;&b6st8>nwNUE>zpaF*6iO^#&O%eV5*y%PY
zBteWBT0_^R>-f2m4IMdjy_cGoEx#maxWfv!w!xquAIx$D0rR(rY|r|dS9PUN@S?ZR
zAu7J!!hY<J@FFt>$jD&rD9o!@MZf2yBv>^MsS1uV9?o7})4xIO{zNYb)89V8oF=>6
zZ%T2ZT^~f6r_Ob@=-yngEje)6Cr%}$>Ga8(TvgS>(;315%5XUfk9$vPwTYkec2Dlg
z9Au@3Qi`hjx?I+Gb2RQVmzFaX2ws}mMYRPe66KLrHdD_&{l0^QHXNRu@m#E4XC(nH
z`|h=_d~<K|IuGGLnb>}_oRsEh$LseY&|&^7SSHFO?^D)OF+dLZ4OXMf@e5pkkv*Cz
zXfaz&oYI4=-}#=Izf0VWKATGSf2owadF#9OfPlJJyV>ju?9i&QI&^XM$5i!yZ!P~{
rpOv2*^ncR^{pZh+e=4Yd&dLv|!k^3cXQ2EHl%KivUzb~@!M^_jdD0iu

diff --git a/examples/screenshots/webgpu_skinning_points.jpg b/examples/screenshots/webgpu_skinning_points.jpg
index 35b06935974bb9787ab26daee5a8712bd8be7e9b..001fe203e0b94e476d12e5c1c76aa979a0fbc77f 100644
GIT binary patch
delta 21
dcmZ2ya#d=>CFV^++8mn`ghCiu{xkf)2>@O+2%`W1

literal 6702
zcmeHLX;c$iwyqE$M&JU<prXxHL2zJ@84+<rhK^i7KnQ_W!5AaZT$x7)RLljDDF}!t
zqYNP;B#{sZAR>dvWtKn~0>~T^1QLd%9_U;B=e-~Hd#n3TyK9|&s_L9t^?m2;v+Fwr
zeg=mD87p&3a{vMX0LaDxfc-!eAOhL6O)s)(qR>q*CMpUQ6&Dj1|5jk|E#fdZOk8}6
z#1{D0P1`t=l-Me{dGY-*5iuxKY%5G0_O1D!a$pT0y9HQ<Jb^-F0TEdUR2Bl(L$Wq7
zihYB&0ra~;M4+N#;;;<}5;7aGq0kMmVjD;|dWLVb0iv>E+x8wkCBEJEGHjo(oX(y2
zms|FqF0GZfYv-uwUitkle5-=uj-9(y52&dhJfx>@U}$7~?9AD7=I1Ret?Vy4I67T&
zcDd@|c@2FXbK{nue?VYRaLB!=`wt#QKf)y>CMDxjo~EW{W#{DP<-aN@d{b6VB2(U0
zQ0v~+zi;@^*woz7+11_C`?-%X!WtbL|1vQ-HO*c4y12B=<FBkj0O)r*8~%Ur*&q%P
z*&K7&CLf4M;6{PUii+($D!%QME$p)I_I)~cw#c21e_2`!->+-Ok-zeL`&I=Ny%AOJ
zCZ%tTevHuF{}ZEs5c&rnkO4?SAsZ70l?BWIVdNz6Lge;VK>cU+ZyEx#0xiBCZ`iPA
zDB6r;;iXg`Vah)`VtVU0pTA(^NkZaCKQ<M2q=01Ic>lGDB=@9pZX~A(1m3g>l$_py
z!2151k^F8z`RBd=xG@M*1_79kuq_f;mj(gRbNoJ5rdXIXCnBzZ0RnY~%U<WLIo$*r
zht1W*$C52ZB`X{~jgCLq|MG?OSo4o9av_mE{OA128Sh6R;0bjDf%c}8%K2TuZ;+E>
zfNzN5oGb^4tUwyAM`F;r(}(ll{Zh6YFRd{X*^(uZ9Mfbp>UTH=(NbYP-bliwghz@l
z521<ZpelRIx6Ouho5}!6|17z($TE~*h>0)>lL=m2^WZm)*dHJS@5xKv!IdJ<v^?62
z3t~n42u91DWGv)=w?ZgE9@?QY4YiPlJ}R?ASs^Fd%>SCZqln@u#E^J7Ls>T~SXww`
z>EnA@8Ctb&hq<v2QVX#-5cu3fjbd#NYiuO!KEiia5kD(nB!Gan(FeU08UrOSbeWkf
zp6@4=+(n!^u<Si#dZY@qINy2A<II{>#HW;9`t;~wzF;`NlU@I%LvY;T+U2LG?^@!_
zQc1LOPju~y!yO~>j11y<B%Zo@eQCkd(x~bEWKX89`NXJJorEdoxM3g{`Of!-&OlXp
zz&0#-#^huCdEEw?YSxTgsfE){_`ahfcXV<UH7_ppa;k(gCd*~B^;MHhudJ0y=P)*5
zaJoMpW42oepJ6;s&Yq0i5vCCuNK_%dWH5prFf8@uqgjr39Gaf)9l@b7Gq{{*WFZ9v
z)OaUqI|K$RQ~9jwMQav4#(nF$6f4}UHbRRlOwONX<k+=%=DQ9aYX0-RWM`AE=B?d_
z_t5H_KcBSso~wsy+FBv1_LZ!s1lYx76jK7+eBONB32T_Yr|6De(|SK7(X`xZfHLoF
z>79$qME?Z@bjGzoK+k%-dJ5KTU|><3Ctx-rEDuo0Dd|X}L$?Bb<odb}2t>(Ot8c{U
z`#<R;J1@l$SXy&I5j8R=>KaA^Q*PIk1YX_02U9+(@|@lBWqEBi@SvIoy~u~RlS-e1
zaMVt9W;*tq&<%1NWSGj`m^+dtc&pquk^6aNC2H)-sfA}TyR4n{D2_JG%cM28>A_b?
zof!9U+FWRVM;kjlwrL$ZJmU`Eczd`%i>YI}rH}BrCx;&65}uKy(>k%etmV)tsb?~$
zTDS4`<h~f4_bPa3HYdhW=(j?JOpFFb?Jvdpy+!tNHC{#6u|yJJI#;n0MOT!J_OXsc
z7P}_!(2a}X)FluY^W>fQs8~f=!n(7mYrVeGb5b19^T7xCqr)7kf3cIxy~sKQ_K!=d
z>ih6EgJIEMF4dLR_#D-~FoGXdNynmn(yDq(;%wADTjpyzR1tkNg=vJaM*lX3{c<3u
zi-H&a%AX%zGuE3RqfQ@=+r}tjq-Sb1ewC}fx($Bi5E6JK_jrq5m&&R?r44237ZNO&
z80jxW3H0?&99X~1599<|zwW8tH#2`3TQIGB<xT1l^t0n%c3#^OmK0uL;lV&c(r<{M
z+L60Rm(Py|2+2#~3UiZ*SEEoi4_86pHpM5aGUUKV4?&?B+Ps6HUr?<xsf}cksQAk2
zONeY6Z=OtUN7I8cG*0bMS#;B3iqdORm4#K(fLb<!nATo2%b1P6XMZ=z-^O98N^pF=
z*U!vdkW($s_w>%~&f&`7+H4iHe0_#vU4nAcJ=lX6<6Yi-!jU@_6LcEW&A*^pp2pD<
z=Z9>bS7{W6>KER8+h8BTuJroS<@9Opun$l7Y7Gb|(6)0GZL2$I2Yp{PX=YT;na`i1
zk>bWH-jYmIo~u}zg<_3XB>r9q0>14}oqXwetELXiAs~Qnv(qf0EeTUQe$)P@hWJq7
zFONYWNiOUb2vnd!z%Km2ZL=qben=)kSF1tb-WmG16je=-Bozt*=T(7-%x@V~4)Q_y
z_Rud<Elj;*oyXG#R*PTKoL(gde^K|0O&a)He`%~)yT!r@xBW9`x8=RaUiR@hFg6}(
z__4pT+w?A$VQd^;<UW?<IB2PLB3J7d;m09(m|C$jHcg5ieL2&nD^1v~)UaCqq_|Zy
zfPd7qJ$6*4KDt!r>7UZc<`-O*@+sFyz2~p-zT(2!CEKkqYVCP6LMt|aKTe+Nn<bD~
z^HD+xZrh`BqGAwE13~orTU~wM3=dy=IR|%I<pJEzxPgD=rbV&N-Gp{Bb!tV>is~ac
z^4g`y-E=}I4|Qa*40~-p7QLpozORhKP94#jWnLS-&_PHW9PC+p&&a52g&WJd22&)x
z<89Q-d}I>>#99yL+Ujc-GRu1@D8JB>xU*aB9E|6Qx+$&Y{VV0^$P*22nQMx|))24d
z^jBv%Q{#!2ijh|o)0lI0P=A>R!&GbCBe^Vx)2>GuxMS+?QfP#+UN5CY4AL{$1nKEf
zoM3g)&el2`5t5U<zFb{N9$QLEZ|8fhX@@EBPx)rIa)#b99^=!b2+ZnKjN&5_jBLU7
z*oLTvJ^OGJ?O>AQ=y|BpqUQar=AH#>wwKUrhK{)3T*8>1|4O23ruz?caw)aS4>M(G
z=U})nTWnz!lT>+TM;TXDkfZ8ioTd8mu1D;EYJK6Jq&M@!F_NcUi5GY)Q^b-#P>VKM
z2N{lZo9Q;{LdcE9?xE*h^sB_ap-Vk4xdsKpe^kagN{P>#Bi!6&r^UI&C1D_N%rLX<
ztIE~o)W-Oy)Ks{mow!oj7vU1IkHIvUB(66(GsW2&0z6~E3%SsDlX_`tG4fufbeJip
zx_@}^UiN@<P{5}?ueoA^zxC6x$NJ$TXJ>v*Cx)-=|CQs5Adl}fWwCmTU9jvof^!~O
zqwRT_If*v&h}74@PiF98N<|q;*Xg?EY>Ah**ErWPi%T7VIeUEW6%8Ku@Y_DxdU3Bq
z-EEP$O1YGCa<mFN+p;OFhGSLW%^XIG&LG>}*nqjlB+oAD(=0iewQkycj9HzQ9P)-D
zL%#b-dF6vq*GkMpluSTXTlz$`?*om!N8pkcD)M@TVb+up%Gh5`TWKJmJhN!TytQn|
zH!Z+$K3`X?R_A~~-|Xzo;lwjs|Jeu!QZfEQe3s?m4@%mvN#+@{XvQ<EVr%<?knY_R
zgz=B-gPFfmkw~^T+il@%uUp|NW#TJ!EgpMbkR;qWbu6!}u&J0|I8F+D*m9ohHM{?U
z#S8L%$G#T{hOyD47~1=^VEEq`bjS=_*HTJ`;Fxm&?m~|vioLWu9Ute-)Cy(t0(<65
zdh`{S1NpA_T@9)7<ci(9X8KGQeS_%sC*!WaCREn`ac9}dR@C_V@i+aDD<AZoB4a_o
z>>Y|s#*^<TusTulDfu(gO9Pc$li~b#HZjlTC>)!5qI6RshWuvb_1V?fdT-*Gr8o!}
z7%Q9}bnTGDt7%lu182KpuqXt|gz4AY;&b6st8>nwNUE>zpaF*6iO^#&O%eV5*y%PY
zBteWBT0_^R>-f2m4IMdjy_cGoEx#maxWfv!w!xquAIx$D0rR(rY|r|dS9PUN@S?ZR
zAu7J!!hY<J@FFt>$jD&rD9o!@MZf2yBv>^MsS1uV9?o7})4xIO{zNYb)89V8oF=>6
zZ%T2ZT^~f6r_Ob@=-yngEje)6Cr%}$>Ga8(TvgS>(;315%5XUfk9$vPwTYkec2Dlg
z9Au@3Qi`hjx?I+Gb2RQVmzFaX2ws}mMYRPe66KLrHdD_&{l0^QHXNRu@m#E4XC(nH
z`|h=_d~<K|IuGGLnb>}_oRsEh$LseY&|&^7SSHFO?^D)OF+dLZ4OXMf@e5pkkv*Cz
zXfaz&oYI4=-}#=Izf0VWKATGSf2owadF#9OfPlJJyV>ju?9i&QI&^XM$5i!yZ!P~{
rpOv2*^ncR^{pZh+e=4Yd&dLv|!k^3cXQ2EHl%KivUzb~@!M^_jdD0iu