| <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> |