<!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 = createBufferWithData(device, { size: data.byteLength, usage: GPUBufferUsage.STORAGE | GPUBufferUsage.MAP_READ }, 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>