1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
|
export const description = `
Stress tests covering behavior around shader entry points.
`;
import { makeTestGroup } from '../../common/framework/test_group.js';
import { range } from '../../common/util/util.js';
import { GPUTest } from '../../webgpu/gpu_test.js';
export const g = makeTestGroup(GPUTest);
const makeCode = (numEntryPoints: number) => {
const kBaseCode = `
struct Buffer { data: u32, };
@group(0) @binding(0) var<storage, read_write> buffer: Buffer;
fn main() { buffer.data = buffer.data + 1u; }
`;
const makeEntryPoint = (i: number) => `
@compute @workgroup_size(1) fn computeMain${i}() { main(); }
`;
return kBaseCode + range(numEntryPoints, makeEntryPoint).join('');
};
g.test('many')
.desc(
`Tests compilation and usage of shaders with a huge number of entry points.
TODO: There may be a normative limit to the number of entry points allowed in
a shader, in which case this would become a validation test instead.`
)
.fn(async t => {
const data = new Uint32Array([0]);
const buffer = t.makeBufferWithContents(data, GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC);
// NOTE: Initial shader compilation time seems to scale exponentially with
// this value in Chrome.
const kNumEntryPoints = 200;
const shader = t.device.createShaderModule({
code: makeCode(kNumEntryPoints),
});
const layout = t.device.createBindGroupLayout({
entries: [
{
binding: 0,
visibility: GPUShaderStage.COMPUTE,
buffer: { type: 'storage' },
},
],
});
const pipelineLayout = t.device.createPipelineLayout({
bindGroupLayouts: [layout],
});
const bindGroup = t.device.createBindGroup({
layout,
entries: [{ binding: 0, resource: { buffer } }],
});
const encoder = t.device.createCommandEncoder();
range(kNumEntryPoints, i => {
const pipeline = t.device.createComputePipeline({
layout: pipelineLayout,
compute: {
module: shader,
entryPoint: `computeMain${i}`,
},
});
const pass = encoder.beginComputePass();
pass.setPipeline(pipeline);
pass.setBindGroup(0, bindGroup);
pass.dispatchWorkgroups(1);
pass.end();
});
t.device.queue.submit([encoder.finish()]);
t.expectGPUBufferValuesEqual(buffer, new Uint32Array([kNumEntryPoints]));
});
|
