blob: e6ae4000c91436754a34e3533314c6d44b2f4801 [file] [log] [blame] [edit]
<script src="../../resources/js-test-pre.js"></script>
<script>
globalThis.testRunner?.waitUntilDone();
const log = globalThis.$vm?.print ?? console.log;
onload = async () => {
let adapter = await navigator.gpu.requestAdapter({});
let device = await adapter.requestDevice({});
device.pushErrorScope('validation');
let code = `
@group(0) @binding(0) var<storage, read_write> buf0: array<vec4u, 16>;
var<workgroup> vb0: bool;
var<workgroup> vb1: bool;
var<workgroup> vb2: bool;
var<workgroup> vb3: bool;
var<workgroup> vu: vec4u;
@compute @workgroup_size(1)
fn c() {
var d0 = u32(vb0);
var d1 = u32(vb1);
var d2 = u32(vb2);
var d3 = u32(vb3);
if d0+d1+d2+d3 < 5 {
buf0[2] = vec4(d0+d1+d2+d3);
}
buf0[d0+d1+d2+d3] = vu;
buf0[0].x = d0;
buf0[0].y = d1;
buf0[0].z = d2;
buf0[0].w = d3;
vb0 = bool(buf0[0][0]);
vb1 = bool(buf0[0][0]);
vb2 = bool(buf0[0][0]);
vb3 = bool(buf0[0][0]);
vu = buf0[0];
}
`;
let module = device.createShaderModule({code});
let bindGroupLayout0 = device.createBindGroupLayout({
entries: [
{binding: 0, buffer: {type: 'storage'}, visibility: GPUShaderStage.COMPUTE},
],
});
let buffer0 = device.createBuffer({
size: 256, usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC,
});
await device.queue.onSubmittedWorkDone();
let bindGroup0 = device.createBindGroup({
layout: bindGroupLayout0, entries: [
{binding: 0, resource: {buffer: buffer0}},
],
});
let pipelineLayout = device.createPipelineLayout({bindGroupLayouts: [bindGroupLayout0]});
let commandEncoder = device.createCommandEncoder();
let computePassEncoder = commandEncoder.beginComputePass({});
let computePipeline = device.createComputePipeline({layout: pipelineLayout, compute: {module}});
computePassEncoder.setPipeline(computePipeline);
computePassEncoder.setBindGroup(0, bindGroup0);
computePassEncoder.dispatchWorkgroups(1);
computePassEncoder.end();
let outputBuffer0 = device.createBuffer({size: buffer0.size, usage: GPUBufferUsage.COPY_DST | GPUBufferUsage.MAP_READ});
commandEncoder.copyBufferToBuffer(buffer0, 0, outputBuffer0, 0, buffer0.size);
device.queue.submit([commandEncoder.finish()]);
await device.queue.onSubmittedWorkDone();
await outputBuffer0.mapAsync(GPUMapMode.READ);
for (x of [...new Uint32Array(outputBuffer0.getMappedRange())]) {
shouldBe('x', '0');
}
outputBuffer0.unmap();
let error = await device.popErrorScope();
if (error) {
log(error.message);
} else {
log('no validation error');
}
globalThis.testRunner?.notifyDone();
};
</script>