|
| 1 | +// irrelevant import so the file becomes a module |
| 2 | +import tgpu from 'typegpu'; |
| 3 | +const t = tgpu; |
| 4 | + |
| 5 | +// setup |
| 6 | +const adapter = await navigator.gpu?.requestAdapter(); |
| 7 | +const device = await adapter?.requestDevice(); |
| 8 | +if (!device) { |
| 9 | + throw new Error('WebGPU is not supported!'); |
| 10 | +} |
| 11 | +const copyModule = device.createShaderModule({ |
| 12 | + label: 'copying compute module', |
| 13 | + code: ` |
| 14 | + struct Item { |
| 15 | + vec: vec3u, |
| 16 | + num: u32, |
| 17 | + } |
| 18 | +
|
| 19 | + @group(0) @binding(0) var<storage, read> sourceBuffer: Item; |
| 20 | + @group(0) @binding(1) var<storage, read_write> targetBuffer: Item; |
| 21 | +
|
| 22 | + @compute @workgroup_size(1) fn computeShader_0(@builtin(global_invocation_id) gid: vec3u){ |
| 23 | + var item = sourceBuffer; |
| 24 | + targetBuffer = item; |
| 25 | + } |
| 26 | + `, |
| 27 | +}); |
| 28 | + |
| 29 | +const pipeline = device.createComputePipeline({ |
| 30 | + label: 'copying compute pipeline', |
| 31 | + layout: 'auto', |
| 32 | + compute: { |
| 33 | + module: copyModule, |
| 34 | + }, |
| 35 | +}); |
| 36 | + |
| 37 | +// work buffer 1 |
| 38 | +const sourceBuffer = device.createBuffer({ |
| 39 | + label: 'source buffer', |
| 40 | + size: 16, |
| 41 | + usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC | |
| 42 | + GPUBufferUsage.COPY_DST, |
| 43 | +}); |
| 44 | + |
| 45 | +// work buffer 2 |
| 46 | +const targetBuffer = device.createBuffer({ |
| 47 | + label: 'target buffer', |
| 48 | + size: 16, |
| 49 | + usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC | |
| 50 | + GPUBufferUsage.COPY_DST, |
| 51 | +}); |
| 52 | + |
| 53 | +// buffer for reading the results |
| 54 | +const resultBuffer = device.createBuffer({ |
| 55 | + label: 'result buffer', |
| 56 | + size: 16, |
| 57 | + usage: GPUBufferUsage.MAP_READ | GPUBufferUsage.COPY_DST, |
| 58 | +}); |
| 59 | + |
| 60 | +const bindGroup = device.createBindGroup({ |
| 61 | + label: 'bind group for work buffers', |
| 62 | + layout: pipeline.getBindGroupLayout(0), |
| 63 | + entries: [ |
| 64 | + { binding: 0, resource: { buffer: sourceBuffer } }, |
| 65 | + { binding: 1, resource: { buffer: targetBuffer } }, |
| 66 | + ], |
| 67 | +}); |
| 68 | + |
| 69 | +// input copying and compute pass |
| 70 | + |
| 71 | +const input = new DataView(new ArrayBuffer(16)); |
| 72 | +input.setUint32(0, 1); |
| 73 | +input.setUint32(4, 3); |
| 74 | +input.setUint32(8, 5); |
| 75 | +input.setUint32(12, 7); |
| 76 | +device.queue.writeBuffer(sourceBuffer, 0, input); |
| 77 | + |
| 78 | +const encoder = device.createCommandEncoder({ |
| 79 | + label: 'copying encoder', |
| 80 | +}); |
| 81 | +const pass = encoder.beginComputePass({ |
| 82 | + label: 'copying compute pass', |
| 83 | +}); |
| 84 | +pass.setPipeline(pipeline); |
| 85 | +pass.setBindGroup(0, bindGroup); |
| 86 | +pass.dispatchWorkgroups(1); |
| 87 | +pass.end(); |
| 88 | + |
| 89 | +encoder.copyBufferToBuffer(targetBuffer, 0, resultBuffer, 0, 16); |
| 90 | +device.queue.submit([encoder.finish()]); |
| 91 | + |
| 92 | +await resultBuffer.mapAsync(GPUMapMode.READ); |
| 93 | +const result = new DataView(resultBuffer.getMappedRange().slice()); |
| 94 | +resultBuffer.unmap(); |
| 95 | + |
| 96 | +const table = document.querySelector<HTMLDivElement>('.result'); |
| 97 | +if (!table) { |
| 98 | + throw new Error('Nowhere to display the results'); |
| 99 | +} |
| 100 | + |
| 101 | +console.log(input, result); |
| 102 | +table.innerText = (input.getUint32(12) === result.getUint32(12)) |
| 103 | + ? 'The bug DOES NOT occur on this device.' |
| 104 | + : 'The bug DOES occur on this device.'; |
0 commit comments