Last active
April 3, 2022 13:03
-
-
Save bellbind/04dbc01c1b80567458f7100d96f47352 to your computer and use it in GitHub Desktop.
[WebGPU] Draw a square with WebGPU for chrome-98/firefox-99
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
// Compute example for WebGPU API: https://www.w3.org/TR/webgpu/ | |
// [Usage] Paste whole codes into Web Comsole of the WebGPU demo page, then output 1024 Float32Array of squares | |
const adapter = await navigator.gpu.requestAdapter(); | |
const device = await adapter.requestDevice(); | |
// WGSL shaders: https://www.w3.org/TR/WGSL/ | |
// NOTE: attribute syntax is changed from [[attr]] to @attr at 2022/01/19; not yet supported on chrome98 | |
// NOTE: from chrome99, [[block]] attribute should be removed | |
const blockAttr = navigator.userAgent?.match(/Chrome\/98/) ? "[[block]]" : ""; | |
const workgroupSize = 64; | |
const computeWgsl = ` | |
${blockAttr} struct IO { | |
values: array<i32>; | |
}; | |
[[binding(0), group(0)]] var<storage, read> input: IO; | |
[[binding(1), group(0)]] var<storage, write> output: IO; | |
[[stage(compute), workgroup_size(${workgroupSize})]] fn square([[builtin(global_invocation_id)]] giid: vec3<u32>) { | |
output.values[giid.x] = input.values[giid.x] * input.values[giid.x]; | |
} | |
`; | |
const computeShader = device.createShaderModule({code: computeWgsl}); | |
// pipeline | |
const pipeline = device.createComputePipeline({ | |
compute: {module: computeShader, entryPoint: "square"}, | |
}); | |
// data | |
const count = 1024; | |
const input = new Int32Array([...Array(count).keys()]); | |
// buffers | |
const inputBuffer = device.createBuffer({size: input.byteLength, usage: GPUBufferUsage.STORAGE, mappedAtCreation: true}); | |
new Int32Array(inputBuffer.getMappedRange()).set(input); | |
inputBuffer.unmap(); | |
const outputBuffer = device.createBuffer({size: input.byteLength, usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC}); | |
const readBuffer = device.createBuffer({size: input.byteLength, usage: GPUBufferUsage.MAP_READ | GPUBufferUsage.COPY_DST}); | |
// bind group | |
const bindGroupLayout = pipeline.getBindGroupLayout(0); | |
const bindGroup = device.createBindGroup({ | |
layout: bindGroupLayout, | |
entries: [ | |
{binding: 0, resource: {buffer: inputBuffer}}, | |
{binding: 1, resource: {buffer: outputBuffer}}, | |
] | |
}); | |
// command encoder | |
const commandEncoder = device.createCommandEncoder(); | |
const passEncoder = commandEncoder.beginComputePass(); | |
passEncoder.setPipeline(pipeline); | |
passEncoder.setBindGroup(0, bindGroup); | |
passEncoder.dispatch(count / workgroupSize); | |
passEncoder.endPass(); //[chrome100] endPass() => end() | |
commandEncoder.copyBufferToBuffer(outputBuffer, 0, readBuffer, 0, input.byteLength); | |
device.queue.submit([commandEncoder.finish()]); | |
// read and copy | |
await readBuffer.mapAsync(GPUMapMode.READ); | |
const output = new Int32Array(readBuffer.getMappedRange().slice()); | |
readBuffer.unmap(); | |
console.log(output); // 1024-result of square 0-1046529 |
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
<!doctype html> | |
<html> | |
<head> | |
<!-- IMPORTANT: The current Chrome requires some origin-trial token in <meta>. | |
To register origins at the last "WebGPU REGISTER" in https://developer.chrome.com/origintrials/ | |
This token is for a Web Origin "https://gist.githack.com" (maybe expired at Mar 31, 2022) | |
It can register localhost origin as "http://localhost:8000" | |
--> | |
<meta http-equiv="origin-trial" | |
content="Akv07qcAop5MFaZYxJtHHjUuM8eV3GpbHkTeuhZo/4wsNjYnQ7GSGJyo7hRVZvpvyjYwilbJ8KbFVchI4O1DpA0AAABQeyJvcmlnaW4iOiJodHRwczovL2dpc3QuZ2l0aGFjay5jb206NDQzIiwiZmVhdHVyZSI6IldlYkdQVSIsImV4cGlyeSI6MTY1MjgzMTk5OX0=" /> | |
<meta http-equiv="origin-trial" | |
content="AkIL+/THBoi1QEsWbX5SOuMpL6+KGAXKrZE5Bz6yHTuijzvKz2MznuLqE+MH4YSqRi/v1fDK/6JyFzgibTTeNAsAAABJeyJvcmlnaW4iOiJodHRwOi8vbG9jYWxob3N0OjgwMDAiLCJmZWF0dXJlIjoiV2ViR1BVIiwiZXhwaXJ5IjoxNjUyODMxOTk5fQ==" /> | |
<script src="./main.js" type="module"></script> | |
<style>@media(prefers-color-scheme: dark){:root {color-scheme: dark;}}</style> | |
<link rel="icon" href="data:image/x-icon;," /> | |
</head> | |
<body> | |
<h1>(Notice: The origin-trial token in this page will be expired at May 15, 2022)</h1> | |
<canvas style="width: 80vmin; height: 80vmin; border: solid;" id="canvas"></canvas> | |
</body> | |
</html> |
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
// Simple example for WebGPU API: https://www.w3.org/TR/webgpu/ | |
const adapter = await navigator.gpu.requestAdapter(); | |
const device = await adapter.requestDevice(); | |
// Buffer of triangle-strip square: 4-(x,y) | |
const square = new Float32Array([-1/2, -1/2, -1/2, +1/2, +1/2, -1/2, +1/2, +1/2,]); | |
const vertexBuffer = device.createBuffer({size: square.byteLength, usage: GPUBufferUsage.VERTEX, mappedAtCreation: true}); | |
new Float32Array(vertexBuffer.getMappedRange()).set(square); | |
vertexBuffer.unmap(); | |
const stride = {arrayStride: 2 * square.BYTES_PER_ELEMENT, attributes: [{shaderLocation: 0, offset: 0, format: "float32x2"}]}; | |
// WGSL shaders: https://www.w3.org/TR/WGSL/ | |
// NOTE: attribute syntax is changed from [[attr]] to @attr at 2022/01/19; not yet supported on chrome98 | |
const vertexWgsl = ` | |
struct Out { | |
[[builtin(position)]] pos: vec4<f32>; | |
[[location(0)]] vert: vec2<f32>; | |
}; | |
[[stage(vertex)]] fn main([[location(0)]] xy: vec2<f32>) -> Out { | |
return Out(vec4<f32>(xy, 0.0, 1.0), xy + 0.5); | |
} | |
`; | |
const vertexShader = device.createShaderModule({code: vertexWgsl}); | |
const fragmentWgsl = ` | |
[[stage(fragment)]] fn main([[builtin(position)]] pos: vec4<f32>, [[location(0)]] vert: vec2<f32>) -> [[location(0)]] vec4<f32> { | |
return vec4<f32>(vert, 0.0, 1.0); | |
} | |
`; | |
const fragmentShader = device.createShaderModule({code: fragmentWgsl}); | |
// gpu config for canvas | |
const canvas = document.getElementById("canvas"); | |
const gpu = canvas.getContext("webgpu"); | |
const format = gpu.getPreferredFormat(adapter); | |
gpu.configure({device, format, size: [canvas.width, canvas.height]}); | |
// pipeline | |
const pipeline = device.createRenderPipeline({ | |
primitive: {topology: "triangle-strip"}, | |
vertex: {module: vertexShader, entryPoint: "main", buffers: [stride]}, | |
fragment: {module: fragmentShader, entryPoint: "main", targets: [{format}]}, | |
}); | |
// render | |
const render = () => { | |
const view = gpu.getCurrentTexture().createView(); | |
//[chrome100] loadValue: {r,g,b,a} => loadOp: "clear", clearValue: {r,g,b,a} (clearValue default is RGBA=0000) | |
const renderPass = {colorAttachments: [{view, loadValue: {r: 0, g: 0, b: 0, a: 0}, storeOp: "store"}]}; | |
const commandEncoder = device.createCommandEncoder(); | |
const passEncoder = commandEncoder.beginRenderPass(renderPass); | |
passEncoder.setPipeline(pipeline); | |
passEncoder.setVertexBuffer(0, vertexBuffer); | |
passEncoder.draw(4, 1); | |
passEncoder.endPass(); //[chrome100] endPass() => end() | |
device.queue.submit([commandEncoder.finish()]); | |
}; | |
(function loop() { | |
render(); | |
requestAnimationFrame(loop); | |
})(); |
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
demo (until Mar 31, 2022): https://gist.githack.com/bellbind/04dbc01c1b80567458f7100d96f47352/raw/index.html
To run on
firefox-nightly
, you should change bothdom.webgpu.enabled
andgfx.webrenderer.all
astrue
atabout:config
panel.To run compute-example.js, paste the whole code into devtool console.
To run compute-example.js on deno-1.19: