| <!DOCTYPE html><!-- webkit-test-runner [ experimental:WebGPUEnabled=true ] --> |
| <meta charset=utf-8> |
| <title>Execute a simple compute shader with Web GPU.</title> |
| <body> |
| <script src="../resources/testharness.js"></script> |
| <script src="../resources/testharnessreport.js"></script> |
| <script src="js/webgpu-functions.js"></script> |
| <script> |
| if (window.testRunner) |
| testRunner.waitUntilDone(); |
| |
| const data = new Uint32Array([2, 3, 4, 5, 6, 7, 8, 9, 10]); |
| |
| const dataBinding = 0; |
| const bindGroupIndex = 0; |
| |
| const shaderCode = ` |
| #include <metal_stdlib> |
| |
| struct Data { |
| device unsigned* numbers [[id(${dataBinding})]]; |
| }; |
| |
| kernel void compute(device Data& data [[buffer(${bindGroupIndex})]], unsigned gid [[thread_position_in_grid]]) |
| { |
| if (gid >= ${data.length}) |
| return; |
| |
| unsigned original = data.numbers[gid]; |
| data.numbers[gid] = original * original; |
| } |
| ` |
| |
| promise_test(async () => { |
| |
| const device = await getBasicDevice(); |
| |
| const shaderModule = device.createShaderModule({ code: shaderCode, isWHLSL: false }); |
| const computeStageDescriptor = { module: shaderModule, entryPoint: "compute" }; |
| const pipeline = device.createComputePipeline({ computeStage: computeStageDescriptor }); |
| |
| const dataBuffer = device.createBuffer({ size: data.byteLength, usage: GPUBufferUsage.STORAGE | GPUBufferUsage.TRANSFER_DST | GPUBufferUsage.MAP_READ }); |
| dataBuffer.setSubData(0, data.buffer); |
| |
| const bgLayoutBinding = { binding: dataBinding, visibility: GPUShaderStageBit.COMPUTE, type: "storage-buffer" }; |
| const bgLayout = device.createBindGroupLayout({ bindings: [bgLayoutBinding] }); |
| |
| const bufferBinding = { buffer: dataBuffer, size: data.byteLength }; |
| const bgBinding = { binding: dataBinding, resource: bufferBinding }; |
| |
| const bindGroupDescriptor = { layout: bgLayout, bindings: [bgBinding] }; |
| const bindGroup = device.createBindGroup(bindGroupDescriptor); |
| |
| const commandEncoder = device.createCommandEncoder(); |
| const passEncoder = commandEncoder.beginComputePass(); |
| |
| passEncoder.setBindGroup(bindGroupIndex, bindGroup); |
| |
| passEncoder.setPipeline(pipeline); |
| |
| // One thread group. |
| passEncoder.dispatch(1, 1, 1); |
| passEncoder.endPass(); |
| |
| device.getQueue().submit([commandEncoder.finish()]); |
| |
| const readDataArrayBuffer = await dataBuffer.mapReadAsync(); |
| assert_not_equals(readDataArrayBuffer, null, "Async read promise resolved successfully"); |
| |
| const readData = new Uint32Array(readDataArrayBuffer); |
| |
| for (var i = 0; i < readData.length; ++i) |
| assert_equals(readData[i], data[i] * data[i], "Data was succesfully squared"); |
| |
| if (window.testRunner) |
| testRunner.notifyDone(); |
| }, "Successfully executed a basic compute pass"); |
| </script> |
| </body> |
