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
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
123
124
125
126
127
128
129
130
131
132
133
134
135
136
137
|
export const description = `
Compound statement execution.
`;
import { makeTestGroup } from '../../../../common/framework/test_group.js';
import { keysOf } from '../../../../common/util/data_tables.js';
import { TypedArrayBufferView } from '../../../../common/util/util.js';
import { GPUTest } from '../../../gpu_test.js';
export const g = makeTestGroup(GPUTest);
/**
* Builds, runs then checks the output of a statement shader test.
*
* @param t The test object
* @param ty The WGSL scalar type to be written
* @param values The expected output values of type ty
* @param wgsl_main The body of the WGSL entry point.
*/
export function runStatementTest(
t: GPUTest,
ty: string,
values: TypedArrayBufferView,
wgsl_main: string
) {
const wgsl = `
struct Outputs {
data : array<${ty}>,
};
var<private> count: u32 = 0;
@group(0) @binding(1) var<storage, read_write> outputs : Outputs;
fn put(value : ${ty}) {
outputs.data[count] = value;
count += 1;
}
@compute @workgroup_size(1)
fn main() {
_ = &outputs;
${wgsl_main}
}
`;
const pipeline = t.device.createComputePipeline({
layout: 'auto',
compute: {
module: t.device.createShaderModule({ code: wgsl }),
entryPoint: 'main',
},
});
const maxOutputValues = 1000;
const outputBuffer = t.device.createBuffer({
size: 4 * (1 + maxOutputValues),
usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC,
});
const bindGroup = t.device.createBindGroup({
layout: pipeline.getBindGroupLayout(0),
entries: [{ binding: 1, resource: { buffer: outputBuffer } }],
});
// Run the shader.
const encoder = t.device.createCommandEncoder();
const pass = encoder.beginComputePass();
pass.setPipeline(pipeline);
pass.setBindGroup(0, bindGroup);
pass.dispatchWorkgroups(1);
pass.end();
t.queue.submit([encoder.finish()]);
t.expectGPUBufferValuesEqual(outputBuffer, values);
}
// Consider a declaration X of identifier 'x' inside a compound statement.
// Check the value of 'x' at various places relative to X:
// a { b; X=c; d; { e; } } f;
const kTests = {
uses: {
// Observe values without conflicting declarations.
src: `let x = 1;
put(x);
{
put(x);
let x = x+1; // The declaration in question
put(x);
{
put(x);
}
put(x);
}
put(x);`,
values: [1, 1, 2, 2, 2, 1],
},
shadowed: {
// Observe values when shadowed
src: `let x = 1;
put(x);
{
put(x);
let x = x+1; // The declaration in question
put(x);
{
let x = x+1; // A shadow
put(x);
}
put(x);
}
put(x);`,
values: [1, 1, 2, 3, 2, 1],
},
gone: {
// The declaration goes out of scope.
src: `{
let x = 2; // The declaration in question
put(x);
}
let x = 1;
put(x);`,
values: [2, 1],
},
} as const;
g.test('decl')
.desc('Tests the value of a declared value in a compound statment.')
.params(u => u.combine('case', keysOf(kTests)))
.fn(t => {
runStatementTest(
t,
'i32',
new Int32Array(kTests[t.params.case].values),
kTests[t.params.case].src
);
});
|
