From 411e08e98e2c468f9d191b5781e90fdcbd9daf15 Mon Sep 17 00:00:00 2001 From: Brandon Jones Date: Tue, 23 Sep 2025 07:26:05 -0700 Subject: [PATCH 1/6] Primitive Picking sample Writes the primitive_index to a secondary render target and uses that to highlight the primitive that the pointer is over on subsequent frames. --- .../computePickPrimitive.wgsl | 14 + .../fragmentForwardRendering.wgsl | 37 ++ .../fragmentPrimitivesDebugView.wgsl | 14 + sample/primitivePicking/index.html | 40 ++ sample/primitivePicking/main.ts | 422 ++++++++++++++++++ sample/primitivePicking/meta.ts | 17 + .../vertexForwardRendering.wgsl | 29 ++ .../primitivePicking/vertexTextureQuad.wgsl | 11 + src/samples.ts | 2 + 9 files changed, 586 insertions(+) create mode 100644 sample/primitivePicking/computePickPrimitive.wgsl create mode 100644 sample/primitivePicking/fragmentForwardRendering.wgsl create mode 100644 sample/primitivePicking/fragmentPrimitivesDebugView.wgsl create mode 100644 sample/primitivePicking/index.html create mode 100644 sample/primitivePicking/main.ts create mode 100644 sample/primitivePicking/meta.ts create mode 100644 sample/primitivePicking/vertexForwardRendering.wgsl create mode 100644 sample/primitivePicking/vertexTextureQuad.wgsl diff --git a/sample/primitivePicking/computePickPrimitive.wgsl b/sample/primitivePicking/computePickPrimitive.wgsl new file mode 100644 index 00000000..6b7c6c88 --- /dev/null +++ b/sample/primitivePicking/computePickPrimitive.wgsl @@ -0,0 +1,14 @@ +struct Frame { + viewProjectionMatrix : mat4x4f, + invViewProjectionMatrix : mat4x4f, + pickCoord: vec2f, + pickedPrimitive: u32, +} +@group(0) @binding(0) var frame : Frame; +@group(0) @binding(1) var primitiveTex: texture_2d; + +@compute @workgroup_size(1) +fn main() { + let texel = vec2u(frame.pickCoord); + frame.pickedPrimitive = textureLoad(primitiveTex, texel, 0).x; +} diff --git a/sample/primitivePicking/fragmentForwardRendering.wgsl b/sample/primitivePicking/fragmentForwardRendering.wgsl new file mode 100644 index 00000000..a7429c86 --- /dev/null +++ b/sample/primitivePicking/fragmentForwardRendering.wgsl @@ -0,0 +1,37 @@ +enable primitive_index; + +struct Frame { + viewProjectionMatrix : mat4x4f, + invViewProjectionMatrix : mat4x4f, + pickCoord: vec2f, + pickedPrimitive: u32, +} +@group(0) @binding(1) var frame : Frame; + +struct PassOutput { + @location(0) color : vec4f, + @location(1) primitive : u32, +} + +@fragment +fn main( + @location(0) fragNormal: vec3f, + @builtin(primitive_index) primIndex: u32 +) -> PassOutput { + let lightDirection = normalize(vec3f(4, 10, 6)); + let light = dot(normalize(fragNormal), lightDirection) * 0.5 + 0.5; + let surfaceColor = vec4f(0.8, 0.8, 0.8, 1.0); + + var output : PassOutput; + + // Highlight the primitive if it's the selected one, otherwise shade normally. + if (primIndex+1 == frame.pickedPrimitive) { + output.color = vec4f(1.0, 1.0, 0.0, 1.0); + } else { + output.color = vec4f(surfaceColor.xyz * light, surfaceColor.a); + } + + // Adding one to each primitive so that 0 can mean "nothing picked" + output.primitive = primIndex+1; + return output; +} diff --git a/sample/primitivePicking/fragmentPrimitivesDebugView.wgsl b/sample/primitivePicking/fragmentPrimitivesDebugView.wgsl new file mode 100644 index 00000000..0b5658c4 --- /dev/null +++ b/sample/primitivePicking/fragmentPrimitivesDebugView.wgsl @@ -0,0 +1,14 @@ +@group(0) @binding(0) var primitiveTex: texture_2d; + +@fragment +fn main( + @builtin(position) coord : vec4f +) -> @location(0) vec4f { + let primitiveIndex = textureLoad(primitiveTex, vec2i(floor(coord.xy)), 0).x; + var result : vec4f; + result.r = f32(primitiveIndex % 8) / 8; + result.g = f32((primitiveIndex / 8) % 8) / 8; + result.b = f32((primitiveIndex / 64) % 8) / 8; + result.a = 1.0; + return result; +} diff --git a/sample/primitivePicking/index.html b/sample/primitivePicking/index.html new file mode 100644 index 00000000..3bb7428a --- /dev/null +++ b/sample/primitivePicking/index.html @@ -0,0 +1,40 @@ + + + + + + webgpu-samples: primitivePicking + + + + + + + + + + + + diff --git a/sample/primitivePicking/main.ts b/sample/primitivePicking/main.ts new file mode 100644 index 00000000..6e27a95b --- /dev/null +++ b/sample/primitivePicking/main.ts @@ -0,0 +1,422 @@ +import { mat4, vec2, vec3, vec4 } from 'wgpu-matrix'; +import { GUI } from 'dat.gui'; +import { mesh } from '../../meshes/teapot'; + +import computePickPrimitive from './computePickPrimitive.wgsl'; +import vertexForwardRendering from './vertexForwardRendering.wgsl'; +import fragmentForwardRendering from './fragmentForwardRendering.wgsl'; +import vertexTextureQuad from './vertexTextureQuad.wgsl'; +import fragmentPrimitivesDebugView from './fragmentPrimitivesDebugView.wgsl'; +import { quitIfWebGPUNotAvailable, quitIfLimitLessThan } from '../util'; + +const canvas = document.querySelector('canvas') as HTMLCanvasElement; +const adapter = await navigator.gpu?.requestAdapter({ + featureLevel: 'compatibility', +}); +const limits: Record = {}; +//@ts-ignore This ignore is needed until primitive-index is added as a feature name to the WebGPU defines. +const features: Array = ['primitive-index']; +quitIfLimitLessThan(adapter, 'maxStorageBuffersInFragmentStage', 1, limits); +const device = await adapter?.requestDevice({ + requiredLimits: limits, + requiredFeatures: features, +}); +quitIfWebGPUNotAvailable(adapter, device); + +const context = canvas.getContext('webgpu') as GPUCanvasContext; + +const devicePixelRatio = window.devicePixelRatio; +canvas.width = canvas.clientWidth * devicePixelRatio; +canvas.height = canvas.clientHeight * devicePixelRatio; +const aspect = canvas.width / canvas.height; +const presentationFormat = navigator.gpu.getPreferredCanvasFormat(); +context.configure({ + device, + format: presentationFormat, +}); + +// Create the model vertex buffer. +const kVertexStride = 8; +const vertexBuffer = device.createBuffer({ + // position: vec3, normal: vec3 + size: mesh.positions.length * kVertexStride * Float32Array.BYTES_PER_ELEMENT, + usage: GPUBufferUsage.VERTEX, + mappedAtCreation: true, +}); +{ + const mapping = new Float32Array(vertexBuffer.getMappedRange()); + for (let i = 0; i < mesh.positions.length; ++i) { + mapping.set(mesh.positions[i], kVertexStride * i); + mapping.set(mesh.normals[i], kVertexStride * i + 3); + } + vertexBuffer.unmap(); +} + +// Create the model index buffer. +const indexCount = mesh.triangles.length * 3; +const indexBuffer = device.createBuffer({ + size: indexCount * Uint16Array.BYTES_PER_ELEMENT, + usage: GPUBufferUsage.INDEX, + mappedAtCreation: true, +}); +{ + const mapping = new Uint16Array(indexBuffer.getMappedRange()); + for (let i = 0; i < mesh.triangles.length; ++i) { + mapping.set(mesh.triangles[i], 3 * i); + } + indexBuffer.unmap(); +} + +// Render targets +const primitiveIndexTexture = device.createTexture({ + size: [canvas.width, canvas.height], + usage: GPUTextureUsage.RENDER_ATTACHMENT | GPUTextureUsage.TEXTURE_BINDING, + format: 'r32uint', +}); +const depthTexture = device.createTexture({ + size: [canvas.width, canvas.height], + format: 'depth24plus', + usage: GPUTextureUsage.RENDER_ATTACHMENT | GPUTextureUsage.TEXTURE_BINDING, +}); + +const vertexBuffers: Iterable = [ + { + arrayStride: Float32Array.BYTES_PER_ELEMENT * 8, + attributes: [ + { + // position + shaderLocation: 0, + offset: 0, + format: 'float32x3', + }, + { + // normal + shaderLocation: 1, + offset: Float32Array.BYTES_PER_ELEMENT * 3, + format: 'float32x3', + }, + { + // uv + shaderLocation: 2, + offset: Float32Array.BYTES_PER_ELEMENT * 6, + format: 'float32x2', + }, + ], + }, +]; + +const primitive: GPUPrimitiveState = { + topology: 'triangle-list', + cullMode: 'back', +}; + +const forwardRenderingPipeline = device.createRenderPipeline({ + layout: 'auto', + vertex: { + module: device.createShaderModule({ + code: vertexForwardRendering, + }), + buffers: vertexBuffers, + }, + fragment: { + module: device.createShaderModule({ + code: fragmentForwardRendering, + }), + targets: [ + // color + { format: presentationFormat }, + // primitive-id + { format: 'r32uint' }, + ], + }, + depthStencil: { + depthWriteEnabled: true, + depthCompare: 'less', + format: 'depth24plus', + }, + primitive, +}); + +const primitiveTextureBindGroupLayout = device.createBindGroupLayout({ + entries: [ + { + binding: 0, + visibility: GPUShaderStage.FRAGMENT, + texture: { + sampleType: 'uint', + }, + }, + ], +}); + +const primitivesDebugViewPipeline = device.createRenderPipeline({ + layout: device.createPipelineLayout({ + bindGroupLayouts: [primitiveTextureBindGroupLayout], + }), + vertex: { + module: device.createShaderModule({ + code: vertexTextureQuad, + }), + }, + fragment: { + module: device.createShaderModule({ + code: fragmentPrimitivesDebugView, + }), + targets: [ + { + format: presentationFormat, + }, + ] + }, + primitive, +}); + +const pickBindGroupLayout = device.createBindGroupLayout({ + entries: [ + { + binding: 0, + visibility: GPUShaderStage.COMPUTE, + buffer: { type: 'storage' } + }, + { + binding: 1, + visibility: GPUShaderStage.COMPUTE, + texture: { + sampleType: 'uint', + }, + }, + ], +}); + +const pickPipeline = device.createComputePipeline({ + layout: device.createPipelineLayout({ + bindGroupLayouts: [ + pickBindGroupLayout, + ], + }), + compute: { + module: device.createShaderModule({ + code: computePickPrimitive, + }), + } +}); + +const forwardRenderPassDescriptor: GPURenderPassDescriptor = { + colorAttachments: [ + { + view: undefined, + + clearValue: [0.0, 0.0, 1.0, 1.0], + loadOp: 'clear', + storeOp: 'store', + }, + { + view: primitiveIndexTexture.createView(), + + clearValue: [0, 0, 0, 1], + loadOp: 'clear', + storeOp: 'store', + }, + ], + depthStencilAttachment: { + view: depthTexture.createView(), + + depthClearValue: 1.0, + depthLoadOp: 'clear', + depthStoreOp: 'store', + }, +}; + +const textureQuadPassDescriptor: GPURenderPassDescriptor = { + colorAttachments: [ + { + // view is acquired and set in render loop. + view: undefined, + + clearValue: [0, 0, 0, 1], + loadOp: 'clear', + storeOp: 'store', + }, + ], +}; + +const settings = { + mode: 'rendering', + rotate: true, +}; +const gui = new GUI(); +gui.add(settings, 'mode', ['rendering', 'primitive indexes']); +gui.add(settings, 'rotate'); + +const modelUniformBuffer = device.createBuffer({ + size: 4 * 16 * 2, // two 4x4 matrix + usage: GPUBufferUsage.UNIFORM | GPUBufferUsage.COPY_DST, +}); + +const cameraUniformBuffer = device.createBuffer({ + size: (4 * 16 * 2) + (4 * 4), // two 4x4 matrix + a u32 + usage: GPUBufferUsage.UNIFORM | GPUBufferUsage.COPY_DST | GPUBufferUsage.STORAGE, +}); + +const sceneUniformBindGroup = device.createBindGroup({ + layout: forwardRenderingPipeline.getBindGroupLayout(0), + entries: [ + { + binding: 0, + resource: { + buffer: modelUniformBuffer, + }, + }, + { + binding: 1, + resource: { + buffer: cameraUniformBuffer, + }, + }, + ], +}); + +const primitiveTextureBindGroup = device.createBindGroup({ + layout: primitiveTextureBindGroupLayout, + entries: [ + { + binding: 0, + resource: primitiveIndexTexture.createView(), + } + ], +}); + +const pickBindGroup = device.createBindGroup({ + layout: pickBindGroupLayout, + entries: [ + { + binding: 0, + resource: cameraUniformBuffer, + }, + { + binding: 1, + resource: primitiveIndexTexture.createView(), + } + ], +}); + +//-------------------- + +// Scene matrices +const eyePosition = vec3.fromValues(0, 12, -25); +const upVector = vec3.fromValues(0, 1, 0); +const origin = vec3.fromValues(0, 0, 0); +const pickCoord = vec2.fromValues(0, 0); + +const projectionMatrix = mat4.perspective((2 * Math.PI) / 5, aspect, 1, 2000.0); + +// Move the model so it's centered. +const modelMatrix = mat4.translation([0, 0, 0]); +device.queue.writeBuffer(modelUniformBuffer, 0, modelMatrix); +const invertTransposeModelMatrix = mat4.invert(modelMatrix); +mat4.transpose(invertTransposeModelMatrix, invertTransposeModelMatrix); +const normalModelData = invertTransposeModelMatrix; +device.queue.writeBuffer( + modelUniformBuffer, + 64, + normalModelData.buffer, + normalModelData.byteOffset, + normalModelData.byteLength +); +// TODO: Write mouse/touch coords + +function onPointerEvent(event: PointerEvent) { + // Only track the primary pointer + if (event.isPrimary) { + let clientRect = (event.target as Element).getBoundingClientRect(); + pickCoord[0] = (event.clientX - clientRect.x) * devicePixelRatio; + pickCoord[1] = (event.clientY - clientRect.y) * devicePixelRatio; + } +} +canvas.addEventListener('pointerenter', onPointerEvent); +canvas.addEventListener('pointermove', onPointerEvent); + +// Rotates the camera around the origin based on time. +let rad = 0; +function getCameraViewProjMatrix() { + if (settings.rotate) { + rad = Math.PI * (Date.now() / 10000); + } + const rotation = mat4.rotateY(mat4.translation(origin), rad); + const rotatedEyePosition = vec3.transformMat4(eyePosition, rotation); + + const viewMatrix = mat4.lookAt(rotatedEyePosition, origin, upVector); + + return mat4.multiply(projectionMatrix, viewMatrix); +} + +function frame() { + const cameraViewProj = getCameraViewProjMatrix(); + device.queue.writeBuffer( + cameraUniformBuffer, + 0, + cameraViewProj.buffer, + cameraViewProj.byteOffset, + cameraViewProj.byteLength + ); + const cameraInvViewProj = mat4.invert(cameraViewProj); + device.queue.writeBuffer( + cameraUniformBuffer, + 64, + cameraInvViewProj.buffer, + cameraInvViewProj.byteOffset, + cameraInvViewProj.byteLength + ); + device.queue.writeBuffer( + cameraUniformBuffer, + 128, + pickCoord.buffer, + pickCoord.byteOffset, + pickCoord.byteLength + ); + + const commandEncoder = device.createCommandEncoder(); + { + // Forward rendering + forwardRenderPassDescriptor.colorAttachments[0].view = context + .getCurrentTexture() + .createView(); + const forwardPass = commandEncoder.beginRenderPass( + forwardRenderPassDescriptor + ); + forwardPass.setPipeline(forwardRenderingPipeline); + forwardPass.setBindGroup(0, sceneUniformBindGroup); + forwardPass.setVertexBuffer(0, vertexBuffer); + forwardPass.setIndexBuffer(indexBuffer, 'uint16'); + forwardPass.drawIndexed(indexCount); + forwardPass.end(); + } + { + if (settings.mode === 'primitive indexes') { + // Primitive Index debug view + // Overwrites the canvas texture with a visualization of the primitive + // index for each primitive + textureQuadPassDescriptor.colorAttachments[0].view = context + .getCurrentTexture() + .createView(); + const debugViewPass = commandEncoder.beginRenderPass( + textureQuadPassDescriptor + ); + debugViewPass.setPipeline(primitivesDebugViewPipeline); + debugViewPass.setBindGroup(0, primitiveTextureBindGroup); + debugViewPass.draw(6); + debugViewPass.end(); + } + } + { + const pickPass = commandEncoder.beginComputePass(); + pickPass.setPipeline(pickPipeline); + pickPass.setBindGroup(0, pickBindGroup); + pickPass.dispatchWorkgroups(1); + pickPass.end(); + } + device.queue.submit([commandEncoder.finish()]); + + requestAnimationFrame(frame); +} +requestAnimationFrame(frame); diff --git a/sample/primitivePicking/meta.ts b/sample/primitivePicking/meta.ts new file mode 100644 index 00000000..133bc345 --- /dev/null +++ b/sample/primitivePicking/meta.ts @@ -0,0 +1,17 @@ +export default { + name: 'Primitive Picking', + description: `This example demonstrates use of the primitive_index builtin. + It is used to render a unique ID for each primitive to a buffer, which is + then read at the current cursor/touch location to determine which primitive + has been selected. That primitive is then highlighted when rendering. + `, + filename: __DIRNAME__, + sources: [ + { path: 'main.ts' }, + { path: 'vertexForwardRendering.wgsl' }, + { path: 'fragmentForwardRendering.wgsl' }, + { path: 'vertexTextureQuad.wgsl' }, + { path: 'fragmentPrimitivesDebugView.wgsl' }, + { path: 'computePickPrimitive.wgsl' }, + ], +}; diff --git a/sample/primitivePicking/vertexForwardRendering.wgsl b/sample/primitivePicking/vertexForwardRendering.wgsl new file mode 100644 index 00000000..28be4750 --- /dev/null +++ b/sample/primitivePicking/vertexForwardRendering.wgsl @@ -0,0 +1,29 @@ +struct Uniforms { + modelMatrix : mat4x4f, + normalModelMatrix : mat4x4f, +} +struct Frame { + viewProjectionMatrix : mat4x4f, + invViewProjectionMatrix : mat4x4f, + pickCoord: vec2u, + pickedPrimitive: u32, +} +@group(0) @binding(0) var uniforms : Uniforms; +@group(0) @binding(1) var frame : Frame; + +struct VertexOutput { + @builtin(position) Position : vec4f, + @location(0) fragNormal: vec3f, // normal in world space +} + +@vertex +fn main( + @location(0) position : vec3f, + @location(1) normal : vec3f, +) -> VertexOutput { + var output : VertexOutput; + let worldPosition = (uniforms.modelMatrix * vec4(position, 1.0)).xyz; + output.Position = frame.viewProjectionMatrix * vec4(worldPosition, 1.0); + output.fragNormal = normalize((uniforms.normalModelMatrix * vec4(normal, 1.0)).xyz); + return output; +} diff --git a/sample/primitivePicking/vertexTextureQuad.wgsl b/sample/primitivePicking/vertexTextureQuad.wgsl new file mode 100644 index 00000000..c1802e7d --- /dev/null +++ b/sample/primitivePicking/vertexTextureQuad.wgsl @@ -0,0 +1,11 @@ +@vertex +fn main( + @builtin(vertex_index) VertexIndex : u32 +) -> @builtin(position) vec4f { + const pos = array( + vec2(-1.0, -1.0), vec2(1.0, -1.0), vec2(-1.0, 1.0), + vec2(-1.0, 1.0), vec2(1.0, -1.0), vec2(1.0, 1.0), + ); + + return vec4f(pos[VertexIndex], 0.0, 1.0); +} diff --git a/src/samples.ts b/src/samples.ts index 8a9ee1ab..5594ed2b 100644 --- a/src/samples.ts +++ b/src/samples.ts @@ -24,6 +24,7 @@ import normalMap from '../sample/normalMap/meta'; import occlusionQuery from '../sample/occlusionQuery/meta'; import particles from '../sample/particles/meta'; import points from '../sample/points/meta'; +import primitivePicking from '../sample/primitivePicking/meta'; import pristineGrid from '../sample/pristineGrid/meta'; import renderBundles from '../sample/renderBundles/meta'; import resizeCanvas from '../sample/resizeCanvas/meta'; @@ -129,6 +130,7 @@ export const pageCategories: PageCategory[] = [ deferredRendering, particles, points, + primitivePicking, imageBlur, generateMipmap, cornell, From 80516918bbc13eebc81146652577d9d320b4e5f8 Mon Sep 17 00:00:00 2001 From: Brandon Jones Date: Tue, 23 Sep 2025 11:53:22 -0700 Subject: [PATCH 2/6] Fix lint errors --- sample/primitivePicking/main.ts | 27 +++++++++++++-------------- sample/primitivePicking/meta.ts | 3 ++- 2 files changed, 15 insertions(+), 15 deletions(-) diff --git a/sample/primitivePicking/main.ts b/sample/primitivePicking/main.ts index 6e27a95b..5a413ca8 100644 --- a/sample/primitivePicking/main.ts +++ b/sample/primitivePicking/main.ts @@ -1,4 +1,4 @@ -import { mat4, vec2, vec3, vec4 } from 'wgpu-matrix'; +import { mat4, vec2, vec3 } from 'wgpu-matrix'; import { GUI } from 'dat.gui'; import { mesh } from '../../meshes/teapot'; @@ -14,7 +14,7 @@ const adapter = await navigator.gpu?.requestAdapter({ featureLevel: 'compatibility', }); const limits: Record = {}; -//@ts-ignore This ignore is needed until primitive-index is added as a feature name to the WebGPU defines. +//@ts-expect-error primitive-index is not yet part of the Typescript definitions. const features: Array = ['primitive-index']; quitIfLimitLessThan(adapter, 'maxStorageBuffersInFragmentStage', 1, limits); const device = await adapter?.requestDevice({ @@ -166,7 +166,7 @@ const primitivesDebugViewPipeline = device.createRenderPipeline({ { format: presentationFormat, }, - ] + ], }, primitive, }); @@ -176,7 +176,7 @@ const pickBindGroupLayout = device.createBindGroupLayout({ { binding: 0, visibility: GPUShaderStage.COMPUTE, - buffer: { type: 'storage' } + buffer: { type: 'storage' }, }, { binding: 1, @@ -190,15 +190,13 @@ const pickBindGroupLayout = device.createBindGroupLayout({ const pickPipeline = device.createComputePipeline({ layout: device.createPipelineLayout({ - bindGroupLayouts: [ - pickBindGroupLayout, - ], + bindGroupLayouts: [pickBindGroupLayout], }), compute: { - module: device.createShaderModule({ + module: device.createShaderModule({ code: computePickPrimitive, }), - } + }, }); const forwardRenderPassDescriptor: GPURenderPassDescriptor = { @@ -254,8 +252,9 @@ const modelUniformBuffer = device.createBuffer({ }); const cameraUniformBuffer = device.createBuffer({ - size: (4 * 16 * 2) + (4 * 4), // two 4x4 matrix + a u32 - usage: GPUBufferUsage.UNIFORM | GPUBufferUsage.COPY_DST | GPUBufferUsage.STORAGE, + size: 4 * 16 * 2 + 4 * 4, // two 4x4 matrix + a u32 + usage: + GPUBufferUsage.UNIFORM | GPUBufferUsage.COPY_DST | GPUBufferUsage.STORAGE, }); const sceneUniformBindGroup = device.createBindGroup({ @@ -282,7 +281,7 @@ const primitiveTextureBindGroup = device.createBindGroup({ { binding: 0, resource: primitiveIndexTexture.createView(), - } + }, ], }); @@ -296,7 +295,7 @@ const pickBindGroup = device.createBindGroup({ { binding: 1, resource: primitiveIndexTexture.createView(), - } + }, ], }); @@ -328,7 +327,7 @@ device.queue.writeBuffer( function onPointerEvent(event: PointerEvent) { // Only track the primary pointer if (event.isPrimary) { - let clientRect = (event.target as Element).getBoundingClientRect(); + const clientRect = (event.target as Element).getBoundingClientRect(); pickCoord[0] = (event.clientX - clientRect.x) * devicePixelRatio; pickCoord[1] = (event.clientY - clientRect.y) * devicePixelRatio; } diff --git a/sample/primitivePicking/meta.ts b/sample/primitivePicking/meta.ts index 133bc345..e2dfece0 100644 --- a/sample/primitivePicking/meta.ts +++ b/sample/primitivePicking/meta.ts @@ -3,7 +3,8 @@ export default { description: `This example demonstrates use of the primitive_index builtin. It is used to render a unique ID for each primitive to a buffer, which is then read at the current cursor/touch location to determine which primitive - has been selected. That primitive is then highlighted when rendering. + has been selected. That primitive is then highlighted when rendering the + next frame. `, filename: __DIRNAME__, sources: [ From b6b9535a865a78a910e0c9fe0e04b846319d9dd8 Mon Sep 17 00:00:00 2001 From: Brandon Jones Date: Tue, 23 Sep 2025 15:13:11 -0700 Subject: [PATCH 3/6] Address feedback from Gregg --- .../computePickPrimitive.wgsl | 2 + .../fragmentForwardRendering.wgsl | 3 +- .../fragmentPrimitivesDebugView.wgsl | 8 +++ sample/primitivePicking/main.ts | 55 ++++++++++--------- sample/util.ts | 13 +++++ 5 files changed, 55 insertions(+), 26 deletions(-) diff --git a/sample/primitivePicking/computePickPrimitive.wgsl b/sample/primitivePicking/computePickPrimitive.wgsl index 6b7c6c88..f9260daa 100644 --- a/sample/primitivePicking/computePickPrimitive.wgsl +++ b/sample/primitivePicking/computePickPrimitive.wgsl @@ -9,6 +9,8 @@ struct Frame { @compute @workgroup_size(1) fn main() { + // Load the primitive index from the picking texture and store it in the + // pickedPrimitive value (exposed to the rendering shaders as a uniform). let texel = vec2u(frame.pickCoord); frame.pickedPrimitive = textureLoad(primitiveTex, texel, 0).x; } diff --git a/sample/primitivePicking/fragmentForwardRendering.wgsl b/sample/primitivePicking/fragmentForwardRendering.wgsl index a7429c86..0fe26ef3 100644 --- a/sample/primitivePicking/fragmentForwardRendering.wgsl +++ b/sample/primitivePicking/fragmentForwardRendering.wgsl @@ -18,6 +18,7 @@ fn main( @location(0) fragNormal: vec3f, @builtin(primitive_index) primIndex: u32 ) -> PassOutput { + // Very simple N-dot-L lighting model let lightDirection = normalize(vec3f(4, 10, 6)); let light = dot(normalize(fragNormal), lightDirection) * 0.5 + 0.5; let surfaceColor = vec4f(0.8, 0.8, 0.8, 1.0); @@ -31,7 +32,7 @@ fn main( output.color = vec4f(surfaceColor.xyz * light, surfaceColor.a); } - // Adding one to each primitive so that 0 can mean "nothing picked" + // Adding one to each primitive index so that 0 can mean "nothing picked" output.primitive = primIndex+1; return output; } diff --git a/sample/primitivePicking/fragmentPrimitivesDebugView.wgsl b/sample/primitivePicking/fragmentPrimitivesDebugView.wgsl index 0b5658c4..9cd156d7 100644 --- a/sample/primitivePicking/fragmentPrimitivesDebugView.wgsl +++ b/sample/primitivePicking/fragmentPrimitivesDebugView.wgsl @@ -4,8 +4,16 @@ fn main( @builtin(position) coord : vec4f ) -> @location(0) vec4f { + // Load the primitive index for this pixel from the picking texture. let primitiveIndex = textureLoad(primitiveTex, vec2i(floor(coord.xy)), 0).x; var result : vec4f; + + // Generate a color for the primitive index. If we only increment the color + // channels by 1 for each primitive index we can show a very large range of + // unique values but it can make the individual primitives hard to distinguish. + // This code steps through 8 distinct values per-channel, which may end up + // repeating some colors for larger meshes but makes the unique primitive + // index values easier to see. result.r = f32(primitiveIndex % 8) / 8; result.g = f32((primitiveIndex / 8) % 8) / 8; result.b = f32((primitiveIndex / 64) % 8) / 8; diff --git a/sample/primitivePicking/main.ts b/sample/primitivePicking/main.ts index 5a413ca8..f0c10bfa 100644 --- a/sample/primitivePicking/main.ts +++ b/sample/primitivePicking/main.ts @@ -7,19 +7,17 @@ import vertexForwardRendering from './vertexForwardRendering.wgsl'; import fragmentForwardRendering from './fragmentForwardRendering.wgsl'; import vertexTextureQuad from './vertexTextureQuad.wgsl'; import fragmentPrimitivesDebugView from './fragmentPrimitivesDebugView.wgsl'; -import { quitIfWebGPUNotAvailable, quitIfLimitLessThan } from '../util'; +import { quitIfWebGPUNotAvailable, quitIfFeaturesNotAvailable } from '../util'; const canvas = document.querySelector('canvas') as HTMLCanvasElement; const adapter = await navigator.gpu?.requestAdapter({ featureLevel: 'compatibility', }); -const limits: Record = {}; //@ts-expect-error primitive-index is not yet part of the Typescript definitions. -const features: Array = ['primitive-index']; -quitIfLimitLessThan(adapter, 'maxStorageBuffersInFragmentStage', 1, limits); +const requiredFeatures: Array = ['primitive-index']; +quitIfFeaturesNotAvailable(adapter, requiredFeatures); const device = await adapter?.requestDevice({ - requiredLimits: limits, - requiredFeatures: features, + requiredFeatures, }); quitIfWebGPUNotAvailable(adapter, device); @@ -68,6 +66,9 @@ const indexBuffer = device.createBuffer({ } // Render targets + +// The primitive index for each triangle will be written out to this texture. +// Using a r32uint texture ensures we can store the full range of primitive indices. const primitiveIndexTexture = device.createTexture({ size: [canvas.width, canvas.height], usage: GPUTextureUsage.RENDER_ATTACHMENT | GPUTextureUsage.TEXTURE_BINDING, @@ -95,19 +96,14 @@ const vertexBuffers: Iterable = [ offset: Float32Array.BYTES_PER_ELEMENT * 3, format: 'float32x3', }, - { - // uv - shaderLocation: 2, - offset: Float32Array.BYTES_PER_ELEMENT * 6, - format: 'float32x2', - }, ], }, ]; const primitive: GPUPrimitiveState = { topology: 'triangle-list', - cullMode: 'back', + // Using `none` because the teapot has gaps that you can see the backfaces through. + cullMode: 'none', }; const forwardRenderingPipeline = device.createRenderPipeline({ @@ -202,6 +198,7 @@ const pickPipeline = device.createComputePipeline({ const forwardRenderPassDescriptor: GPURenderPassDescriptor = { colorAttachments: [ { + // view is acquired and set in render loop. view: undefined, clearValue: [0.0, 0.0, 1.0, 1.0], @@ -211,7 +208,6 @@ const forwardRenderPassDescriptor: GPURenderPassDescriptor = { { view: primitiveIndexTexture.createView(), - clearValue: [0, 0, 0, 1], loadOp: 'clear', storeOp: 'store', }, @@ -246,13 +242,16 @@ const gui = new GUI(); gui.add(settings, 'mode', ['rendering', 'primitive indexes']); gui.add(settings, 'rotate'); +const MatrixSizeBytes = Float32Array.BYTES_PER_ELEMENT * 16; +const PickUniformsSizeBytes = Float32Array.BYTES_PER_ELEMENT * 4; + const modelUniformBuffer = device.createBuffer({ - size: 4 * 16 * 2, // two 4x4 matrix + size: MatrixSizeBytes * 2, // two 4x4 matrix usage: GPUBufferUsage.UNIFORM | GPUBufferUsage.COPY_DST, }); -const cameraUniformBuffer = device.createBuffer({ - size: 4 * 16 * 2 + 4 * 4, // two 4x4 matrix + a u32 +const frameUniformBuffer = device.createBuffer({ + size: MatrixSizeBytes * 2 + PickUniformsSizeBytes, // two 4x4 matrix + a vec4's worth of picking uniforms usage: GPUBufferUsage.UNIFORM | GPUBufferUsage.COPY_DST | GPUBufferUsage.STORAGE, }); @@ -269,7 +268,7 @@ const sceneUniformBindGroup = device.createBindGroup({ { binding: 1, resource: { - buffer: cameraUniformBuffer, + buffer: frameUniformBuffer, }, }, ], @@ -290,7 +289,7 @@ const pickBindGroup = device.createBindGroup({ entries: [ { binding: 0, - resource: cameraUniformBuffer, + resource: frameUniformBuffer, }, { binding: 1, @@ -305,7 +304,6 @@ const pickBindGroup = device.createBindGroup({ const eyePosition = vec3.fromValues(0, 12, -25); const upVector = vec3.fromValues(0, 1, 0); const origin = vec3.fromValues(0, 0, 0); -const pickCoord = vec2.fromValues(0, 0); const projectionMatrix = mat4.perspective((2 * Math.PI) / 5, aspect, 1, 2000.0); @@ -322,12 +320,14 @@ device.queue.writeBuffer( normalModelData.byteOffset, normalModelData.byteLength ); -// TODO: Write mouse/touch coords +// Pointer tracking +const pickCoord = vec2.fromValues(0, 0); function onPointerEvent(event: PointerEvent) { // Only track the primary pointer if (event.isPrimary) { const clientRect = (event.target as Element).getBoundingClientRect(); + // Get the pixel offset from the top-left of the canvas element. pickCoord[0] = (event.clientX - clientRect.x) * devicePixelRatio; pickCoord[1] = (event.clientY - clientRect.y) * devicePixelRatio; } @@ -352,7 +352,7 @@ function getCameraViewProjMatrix() { function frame() { const cameraViewProj = getCameraViewProjMatrix(); device.queue.writeBuffer( - cameraUniformBuffer, + frameUniformBuffer, 0, cameraViewProj.buffer, cameraViewProj.byteOffset, @@ -360,14 +360,14 @@ function frame() { ); const cameraInvViewProj = mat4.invert(cameraViewProj); device.queue.writeBuffer( - cameraUniformBuffer, + frameUniformBuffer, 64, cameraInvViewProj.buffer, cameraInvViewProj.byteOffset, cameraInvViewProj.byteLength ); device.queue.writeBuffer( - cameraUniformBuffer, + frameUniformBuffer, 128, pickCoord.buffer, pickCoord.byteOffset, @@ -376,7 +376,7 @@ function frame() { const commandEncoder = device.createCommandEncoder(); { - // Forward rendering + // Forward rendering pass forwardRenderPassDescriptor.colorAttachments[0].view = context .getCurrentTexture() .createView(); @@ -408,6 +408,11 @@ function frame() { } } { + // Picking pass. Executes a single instance of a compue shader that loads + // the primitive index at the pointer coordinates from the primitive index + // texture written in the forward pass. The selected primitive index is + // saved in the frameUniformBuffer and used for highlighting on the next + // render. This means that the highlighted primitive is always a frame behind. const pickPass = commandEncoder.beginComputePass(); pickPass.setPipeline(pickPipeline); pickPass.setBindGroup(0, pickBindGroup); diff --git a/sample/util.ts b/sample/util.ts index 2628ba48..c249f350 100644 --- a/sample/util.ts +++ b/sample/util.ts @@ -40,6 +40,19 @@ export function quitIfLimitLessThan( } } +export function quitIfFeaturesNotAvailable( + adapter: GPUAdapter, + requiredFeatures: GPUFeatureName[] +) { + for (const feature of requiredFeatures) { + if (!adapter.features.has(feature)) { + fail( + `This sample requries the '${feature}' feature, which is not supported by this system.` + ); + } + } +} + /** * Shows an error dialog if getting a adapter or device wasn't successful, * or if/when the device is lost or has an uncaptured error. From 5f35c34b52a0b2fcfd119ac3502db167f61bfab6 Mon Sep 17 00:00:00 2001 From: Brandon Jones Date: Tue, 23 Sep 2025 15:15:59 -0700 Subject: [PATCH 4/6] More consistent WGSL formatting --- sample/primitivePicking/computePickPrimitive.wgsl | 6 +++--- sample/primitivePicking/fragmentForwardRendering.wgsl | 8 ++++---- sample/primitivePicking/vertexForwardRendering.wgsl | 6 +++--- 3 files changed, 10 insertions(+), 10 deletions(-) diff --git a/sample/primitivePicking/computePickPrimitive.wgsl b/sample/primitivePicking/computePickPrimitive.wgsl index f9260daa..555179f1 100644 --- a/sample/primitivePicking/computePickPrimitive.wgsl +++ b/sample/primitivePicking/computePickPrimitive.wgsl @@ -1,11 +1,11 @@ struct Frame { viewProjectionMatrix : mat4x4f, invViewProjectionMatrix : mat4x4f, - pickCoord: vec2f, - pickedPrimitive: u32, + pickCoord : vec2f, + pickedPrimitive : u32, } @group(0) @binding(0) var frame : Frame; -@group(0) @binding(1) var primitiveTex: texture_2d; +@group(0) @binding(1) var primitiveTex : texture_2d; @compute @workgroup_size(1) fn main() { diff --git a/sample/primitivePicking/fragmentForwardRendering.wgsl b/sample/primitivePicking/fragmentForwardRendering.wgsl index 0fe26ef3..873dda53 100644 --- a/sample/primitivePicking/fragmentForwardRendering.wgsl +++ b/sample/primitivePicking/fragmentForwardRendering.wgsl @@ -3,8 +3,8 @@ enable primitive_index; struct Frame { viewProjectionMatrix : mat4x4f, invViewProjectionMatrix : mat4x4f, - pickCoord: vec2f, - pickedPrimitive: u32, + pickCoord : vec2f, + pickedPrimitive : u32, } @group(0) @binding(1) var frame : Frame; @@ -15,8 +15,8 @@ struct PassOutput { @fragment fn main( - @location(0) fragNormal: vec3f, - @builtin(primitive_index) primIndex: u32 + @location(0) fragNormal : vec3f, + @builtin(primitive_index) primIndex : u32 ) -> PassOutput { // Very simple N-dot-L lighting model let lightDirection = normalize(vec3f(4, 10, 6)); diff --git a/sample/primitivePicking/vertexForwardRendering.wgsl b/sample/primitivePicking/vertexForwardRendering.wgsl index 28be4750..34c8cc88 100644 --- a/sample/primitivePicking/vertexForwardRendering.wgsl +++ b/sample/primitivePicking/vertexForwardRendering.wgsl @@ -5,15 +5,15 @@ struct Uniforms { struct Frame { viewProjectionMatrix : mat4x4f, invViewProjectionMatrix : mat4x4f, - pickCoord: vec2u, - pickedPrimitive: u32, + pickCoord : vec2u, + pickedPrimitive : u32, } @group(0) @binding(0) var uniforms : Uniforms; @group(0) @binding(1) var frame : Frame; struct VertexOutput { @builtin(position) Position : vec4f, - @location(0) fragNormal: vec3f, // normal in world space + @location(0) fragNormal : vec3f, // normal in world space } @vertex From 22f1503d8cf44046400d86fab7ea7524758e4af4 Mon Sep 17 00:00:00 2001 From: Brandon Jones Date: Wed, 24 Sep 2025 08:24:12 -0700 Subject: [PATCH 5/6] Address feedback from Francois --- package-lock.json | 9 +++++---- sample/primitivePicking/main.ts | 19 ++++++++++--------- sample/primitivePicking/meta.ts | 2 +- sample/util.ts | 13 ++++++++++--- 4 files changed, 26 insertions(+), 17 deletions(-) diff --git a/package-lock.json b/package-lock.json index b8d1ba12..ec6ad6ee 100644 --- a/package-lock.json +++ b/package-lock.json @@ -1103,10 +1103,11 @@ "dev": true }, "node_modules/@webgpu/types": { - "version": "0.1.61", - "resolved": "https://registry.npmjs.org/@webgpu/types/-/types-0.1.61.tgz", - "integrity": "sha512-w2HbBvH+qO19SB5pJOJFKs533CdZqxl3fcGonqL321VHkW7W/iBo6H8bjDy6pr/+pbMwIu5dnuaAxH7NxBqUrQ==", - "dev": true + "version": "0.1.65", + "resolved": "https://registry.npmjs.org/@webgpu/types/-/types-0.1.65.tgz", + "integrity": "sha512-cYrHab4d6wuVvDW5tdsfI6/o6vcLMDe6w2Citd1oS51Xxu2ycLCnVo4fqwujfKWijrZMInTJIKcXxteoy21nVA==", + "dev": true, + "license": "BSD-3-Clause" }, "node_modules/accepts": { "version": "1.3.8", diff --git a/sample/primitivePicking/main.ts b/sample/primitivePicking/main.ts index f0c10bfa..3b816d3b 100644 --- a/sample/primitivePicking/main.ts +++ b/sample/primitivePicking/main.ts @@ -13,10 +13,11 @@ const canvas = document.querySelector('canvas') as HTMLCanvasElement; const adapter = await navigator.gpu?.requestAdapter({ featureLevel: 'compatibility', }); -//@ts-expect-error primitive-index is not yet part of the Typescript definitions. + const requiredFeatures: Array = ['primitive-index']; quitIfFeaturesNotAvailable(adapter, requiredFeatures); -const device = await adapter?.requestDevice({ + +const device = await adapter.requestDevice({ requiredFeatures, }); quitIfWebGPUNotAvailable(adapter, device); @@ -34,7 +35,7 @@ context.configure({ }); // Create the model vertex buffer. -const kVertexStride = 8; +const kVertexStride = 6; const vertexBuffer = device.createBuffer({ // position: vec3, normal: vec3 size: mesh.positions.length * kVertexStride * Float32Array.BYTES_PER_ELEMENT, @@ -82,7 +83,7 @@ const depthTexture = device.createTexture({ const vertexBuffers: Iterable = [ { - arrayStride: Float32Array.BYTES_PER_ELEMENT * 8, + arrayStride: Float32Array.BYTES_PER_ELEMENT * kVertexStride, attributes: [ { // position @@ -242,16 +243,16 @@ const gui = new GUI(); gui.add(settings, 'mode', ['rendering', 'primitive indexes']); gui.add(settings, 'rotate'); -const MatrixSizeBytes = Float32Array.BYTES_PER_ELEMENT * 16; -const PickUniformsSizeBytes = Float32Array.BYTES_PER_ELEMENT * 4; +const kMatrixSizeBytes = Float32Array.BYTES_PER_ELEMENT * 16; +const kPickUniformsSizeBytes = Float32Array.BYTES_PER_ELEMENT * 4; const modelUniformBuffer = device.createBuffer({ - size: MatrixSizeBytes * 2, // two 4x4 matrix + size: kMatrixSizeBytes * 2, // two 4x4 matrix usage: GPUBufferUsage.UNIFORM | GPUBufferUsage.COPY_DST, }); const frameUniformBuffer = device.createBuffer({ - size: MatrixSizeBytes * 2 + PickUniformsSizeBytes, // two 4x4 matrix + a vec4's worth of picking uniforms + size: kMatrixSizeBytes * 2 + kPickUniformsSizeBytes, // two 4x4 matrix + a vec4's worth of picking uniforms usage: GPUBufferUsage.UNIFORM | GPUBufferUsage.COPY_DST | GPUBufferUsage.STORAGE, }); @@ -408,7 +409,7 @@ function frame() { } } { - // Picking pass. Executes a single instance of a compue shader that loads + // Picking pass. Executes a single instance of a compute shader that loads // the primitive index at the pointer coordinates from the primitive index // texture written in the forward pass. The selected primitive index is // saved in the frameUniformBuffer and used for highlighting on the next diff --git a/sample/primitivePicking/meta.ts b/sample/primitivePicking/meta.ts index e2dfece0..34c7fe9d 100644 --- a/sample/primitivePicking/meta.ts +++ b/sample/primitivePicking/meta.ts @@ -1,6 +1,6 @@ export default { name: 'Primitive Picking', - description: `This example demonstrates use of the primitive_index builtin. + description: `This example demonstrates use of the primitive_index WGSL builtin. It is used to render a unique ID for each primitive to a buffer, which is then read at the current cursor/touch location to determine which primitive has been selected. That primitive is then highlighted when rendering the diff --git a/sample/util.ts b/sample/util.ts index c249f350..5a392278 100644 --- a/sample/util.ts +++ b/sample/util.ts @@ -40,15 +40,22 @@ export function quitIfLimitLessThan( } } +/** + * Shows an error dialog if getting an adapter wasn't successful or the adapter + * does not support the given list of features. + */ export function quitIfFeaturesNotAvailable( - adapter: GPUAdapter, + adapter: GPUAdapter | null, requiredFeatures: GPUFeatureName[] -) { +): asserts adapter { + quitIfAdapterNotAvailable(adapter); + for (const feature of requiredFeatures) { if (!adapter.features.has(feature)) { fail( - `This sample requries the '${feature}' feature, which is not supported by this system.` + `This sample requires the '${feature}' feature, which is not supported by this system.` ); + return; } } } From 2d82e0f6d5a0359d8489d2f4f073da1c93c0a3f5 Mon Sep 17 00:00:00 2001 From: Brandon Jones Date: Thu, 25 Sep 2025 09:39:05 -0700 Subject: [PATCH 6/6] Update sample/primitivePicking/main.ts MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Co-authored-by: François Beaufort --- sample/primitivePicking/main.ts | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sample/primitivePicking/main.ts b/sample/primitivePicking/main.ts index 3b816d3b..8966f24f 100644 --- a/sample/primitivePicking/main.ts +++ b/sample/primitivePicking/main.ts @@ -14,7 +14,7 @@ const adapter = await navigator.gpu?.requestAdapter({ featureLevel: 'compatibility', }); -const requiredFeatures: Array = ['primitive-index']; +const requiredFeatures: GPUFeatureName[] = ['primitive-index']; quitIfFeaturesNotAvailable(adapter, requiredFeatures); const device = await adapter.requestDevice({