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/computePickPrimitive.wgsl b/sample/primitivePicking/computePickPrimitive.wgsl new file mode 100644 index 00000000..555179f1 --- /dev/null +++ b/sample/primitivePicking/computePickPrimitive.wgsl @@ -0,0 +1,16 @@ +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() { + // 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 new file mode 100644 index 00000000..873dda53 --- /dev/null +++ b/sample/primitivePicking/fragmentForwardRendering.wgsl @@ -0,0 +1,38 @@ +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 { + // 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); + + 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 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 new file mode 100644 index 00000000..9cd156d7 --- /dev/null +++ b/sample/primitivePicking/fragmentPrimitivesDebugView.wgsl @@ -0,0 +1,22 @@ +@group(0) @binding(0) var primitiveTex: texture_2d; + +@fragment +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; + 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..8966f24f --- /dev/null +++ b/sample/primitivePicking/main.ts @@ -0,0 +1,427 @@ +import { mat4, vec2, vec3 } 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, quitIfFeaturesNotAvailable } from '../util'; + +const canvas = document.querySelector('canvas') as HTMLCanvasElement; +const adapter = await navigator.gpu?.requestAdapter({ + featureLevel: 'compatibility', +}); + +const requiredFeatures: GPUFeatureName[] = ['primitive-index']; +quitIfFeaturesNotAvailable(adapter, requiredFeatures); + +const device = await adapter.requestDevice({ + requiredFeatures, +}); +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 = 6; +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 + +// 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, + 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 * kVertexStride, + attributes: [ + { + // position + shaderLocation: 0, + offset: 0, + format: 'float32x3', + }, + { + // normal + shaderLocation: 1, + offset: Float32Array.BYTES_PER_ELEMENT * 3, + format: 'float32x3', + }, + ], + }, +]; + +const primitive: GPUPrimitiveState = { + topology: 'triangle-list', + // Using `none` because the teapot has gaps that you can see the backfaces through. + cullMode: 'none', +}; + +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 is acquired and set in render loop. + view: undefined, + + clearValue: [0.0, 0.0, 1.0, 1.0], + loadOp: 'clear', + storeOp: 'store', + }, + { + view: primitiveIndexTexture.createView(), + + 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 kMatrixSizeBytes = Float32Array.BYTES_PER_ELEMENT * 16; +const kPickUniformsSizeBytes = Float32Array.BYTES_PER_ELEMENT * 4; + +const modelUniformBuffer = device.createBuffer({ + size: kMatrixSizeBytes * 2, // two 4x4 matrix + usage: GPUBufferUsage.UNIFORM | GPUBufferUsage.COPY_DST, +}); + +const frameUniformBuffer = device.createBuffer({ + size: kMatrixSizeBytes * 2 + kPickUniformsSizeBytes, // two 4x4 matrix + a vec4's worth of picking uniforms + 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: frameUniformBuffer, + }, + }, + ], +}); + +const primitiveTextureBindGroup = device.createBindGroup({ + layout: primitiveTextureBindGroupLayout, + entries: [ + { + binding: 0, + resource: primitiveIndexTexture.createView(), + }, + ], +}); + +const pickBindGroup = device.createBindGroup({ + layout: pickBindGroupLayout, + entries: [ + { + binding: 0, + resource: frameUniformBuffer, + }, + { + 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 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 +); + +// 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; + } +} +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( + frameUniformBuffer, + 0, + cameraViewProj.buffer, + cameraViewProj.byteOffset, + cameraViewProj.byteLength + ); + const cameraInvViewProj = mat4.invert(cameraViewProj); + device.queue.writeBuffer( + frameUniformBuffer, + 64, + cameraInvViewProj.buffer, + cameraInvViewProj.byteOffset, + cameraInvViewProj.byteLength + ); + device.queue.writeBuffer( + frameUniformBuffer, + 128, + pickCoord.buffer, + pickCoord.byteOffset, + pickCoord.byteLength + ); + + const commandEncoder = device.createCommandEncoder(); + { + // Forward rendering pass + 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(); + } + } + { + // 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 + // render. This means that the highlighted primitive is always a frame behind. + 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..34c7fe9d --- /dev/null +++ b/sample/primitivePicking/meta.ts @@ -0,0 +1,18 @@ +export default { + name: 'Primitive Picking', + 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 + next frame. + `, + 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..34c8cc88 --- /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/sample/util.ts b/sample/util.ts index 2628ba48..5a392278 100644 --- a/sample/util.ts +++ b/sample/util.ts @@ -40,6 +40,26 @@ 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 | null, + requiredFeatures: GPUFeatureName[] +): asserts adapter { + quitIfAdapterNotAvailable(adapter); + + for (const feature of requiredFeatures) { + if (!adapter.features.has(feature)) { + fail( + `This sample requires the '${feature}' feature, which is not supported by this system.` + ); + return; + } + } +} + /** * 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. 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,