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
|
export const description = `
Validation tests for the workgroupUniformLoad() builtin.
`;
import { makeTestGroup } from '../../../../../../common/framework/test_group.js';
import { keysOf } from '../../../../../../common/util/data_tables.js';
import { ShaderValidationTest } from '../../../shader_validation_test.js';
export const g = makeTestGroup(ShaderValidationTest);
const kEntryPoints = {
none: { supportsBarrier: true, code: `` },
compute: {
supportsBarrier: true,
code: `@compute @workgroup_size(1)
fn main() {
foo();
}`,
},
vertex: {
supportsBarrier: false,
code: `@vertex
fn main() -> @builtin(position) vec4f {
foo();
return vec4f();
}`,
},
fragment: {
supportsBarrier: false,
code: `@fragment
fn main() {
foo();
}`,
},
compute_and_fragment: {
supportsBarrier: false,
code: `@compute @workgroup_size(1)
fn main1() {
foo();
}
@fragment
fn main2() {
foo();
}
`,
},
fragment_without_call: {
supportsBarrier: true,
code: `@fragment
fn main() {
}
`,
},
};
g.test('only_in_compute')
.specURL('https://www.w3.org/TR/WGSL/#sync-builtin-functions')
.desc(
`
Synchronization functions must only be used in the compute shader stage.
`
)
.params(u =>
u
.combine('entry_point', keysOf(kEntryPoints))
.combine('call', ['bar()', 'workgroupUniformLoad(&wgvar)'])
)
.fn(t => {
const config = kEntryPoints[t.params.entry_point];
const code = `
${config.code}
var<workgroup> wgvar : u32;
fn bar() -> u32 {
return 0;
}
fn foo() {
_ = ${t.params.call};
}`;
t.expectCompileResult(t.params.call === 'bar()' || config.supportsBarrier, code);
});
// A list of types that contains atomics, with a single control case.
const kAtomicTypes: string[] = [
'bool', // control case
'atomic<i32>',
'atomic<u32>',
'array<atomic<i32>, 4>',
'AtomicStruct',
];
g.test('no_atomics')
.desc(
`
The argument passed to workgroupUniformLoad cannot contain any atomic types.
NOTE: Various other valid types are tested via execution tests, so we only check for invalid types here.
`
)
.params(u =>
u.combine('type', kAtomicTypes).combine('call', ['bar()', 'workgroupUniformLoad(&wgvar)'])
)
.fn(t => {
const code = `
struct AtomicStruct {
a : atomic<u32>
}
var<workgroup> wgvar : ${t.params.type};
fn bar() -> bool {
return true;
}
fn foo() {
_ = ${t.params.call};
}`;
t.expectCompileResult(t.params.type === 'bool' || t.params.call === 'bar()', code);
});
|
