Add insertBits edge case regression test (#4664)
This reproduces a GPU hang on Mac for the Dawn implementation (see crbug.com/506180954).
This commit is contained in:
parent
12200d24a6
commit
d7791c8417
2 changed files with 44 additions and 0 deletions
|
|
@ -1491,6 +1491,7 @@
|
|||
"webgpu:shader,execution,expression,call,builtin,fwidth:f32:*": { "subcaseMS": 29.807 },
|
||||
"webgpu:shader,execution,expression,call,builtin,fwidthCoarse:f32:*": { "subcaseMS": 17.110 },
|
||||
"webgpu:shader,execution,expression,call,builtin,fwidthFine:f32:*": { "subcaseMS": 16.906 },
|
||||
"webgpu:shader,execution,expression,call,builtin,insertBits:array_index:*": { "subcaseMS": 20.682 },
|
||||
"webgpu:shader,execution,expression,call,builtin,insertBits:integer:*": { "subcaseMS": 9.569 },
|
||||
"webgpu:shader,execution,expression,call,builtin,inversesqrt:abstract_float:*": { "subcaseMS": 19408.045 },
|
||||
"webgpu:shader,execution,expression,call,builtin,inversesqrt:f16:*": { "subcaseMS": 21.411 },
|
||||
|
|
|
|||
|
|
@ -378,3 +378,46 @@ g.test('integer')
|
|||
|
||||
await run(t, builtin('insertBits'), [T, T, Type.u32, Type.u32], T, cfg, cases);
|
||||
});
|
||||
|
||||
g.test('array_index')
|
||||
.desc('Test insertBits offset edge case where the result is used as an array index.')
|
||||
.fn(t => {
|
||||
const wgsl = `
|
||||
@group(0) @binding(0) var<storage, read_write> res: array<u32, 8>;
|
||||
|
||||
@compute @workgroup_size(1)
|
||||
fn main() {
|
||||
let count = 0u;
|
||||
res[insertBits(5u, 1u, 32u, count)] = 42u;
|
||||
}
|
||||
`;
|
||||
|
||||
const pipeline = t.device.createComputePipeline({
|
||||
layout: 'auto',
|
||||
compute: {
|
||||
module: t.device.createShaderModule({ code: wgsl }),
|
||||
entryPoint: 'main',
|
||||
},
|
||||
});
|
||||
|
||||
const resBuffer = t.createBufferTracked({
|
||||
size: 8 * 4,
|
||||
usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC,
|
||||
});
|
||||
|
||||
const group = t.device.createBindGroup({
|
||||
layout: pipeline.getBindGroupLayout(0),
|
||||
entries: [{ binding: 0, resource: { buffer: resBuffer } }],
|
||||
});
|
||||
|
||||
const encoder = t.device.createCommandEncoder();
|
||||
const pass = encoder.beginComputePass();
|
||||
pass.setPipeline(pipeline);
|
||||
pass.setBindGroup(0, group);
|
||||
pass.dispatchWorkgroups(1);
|
||||
pass.end();
|
||||
|
||||
t.queue.submit([encoder.finish()]);
|
||||
|
||||
t.expectGPUBufferValuesEqual(resBuffer, new Uint32Array([0, 0, 0, 0, 0, 42, 0, 0]));
|
||||
});
|
||||
|
|
|
|||
Loading…
Reference in a new issue