Skip to content

Commit 9589788

Browse files
mwyrzykowskimwyrzykowski
committed
WGSL scoping bug hides a duplicate @binding, undersizing the argument buffer and forging a shader device pointer
https://bugs.webkit.org/show_bug.cgi?id=323904 rdar://186926179 Reviewed by Gerald Squelart. Visit the declaration before introducing its name in both visitors, and turn the duplicate-binding ASSERT in determineUsedGlobals into a hard error so a future divergence between the validator and the rewriter cannot produce a shader whose argument buffer disagrees with its layout in a release build. Tests: sar3/fast/webgpu/regression/repro_186926179.html sar3/fast/webgpu/regression/shader-module-declaration-scope.html * LayoutTests/sar3/fast/webgpu/regression/repro_186926179-expected.txt: Added. * LayoutTests/sar3/fast/webgpu/regression/repro_186926179.html: Added. * LayoutTests/sar3/fast/webgpu/regression/shader-module-declaration-scope-expected.txt: Added. * LayoutTests/sar3/fast/webgpu/regression/shader-module-declaration-scope.html: Added. * Source/WebGPU/WGSL/GlobalSorting.cpp: (WGSL::GraphBuilder::visit): * Source/WebGPU/WGSL/GlobalVariableRewriter.cpp: (WGSL::RewriteGlobalVariables::determineUsedGlobals): * Source/WebGPU/WGSL/IOValidator.cpp: (WGSL::IOValidator::visit): * Tools/TestWebKitAPI/Tests/WGSL/TypeCheckingTests.mm: (TestWebKitAPI::expectStaticCheckErrorContaining): (TestWebKitAPI::TEST): Originally-landed-as: 316606.692@safari-7625-branch (7e32e6e). rdar://191050941 Canonical link: https://commits.webkit.org/325176@main
1 parent 66be4cf commit 9589788

8 files changed

Lines changed: 324 additions & 5 deletions

File tree

‎LayoutTests/sar3/fast/webgpu/regression/repro_186926179-expected.txt‎

Lines changed: 7 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,7 @@
1+
PASS the duplicate resource binding was reported
2+
PASS dispatching the invalid module did not crash the GPU process
3+
PASS an entry point using only one of the two bindings still runs
4+
PASS successfullyParsed is true
5+
6+
TEST COMPLETE
7+

‎LayoutTests/sar3/fast/webgpu/regression/repro_186926179.html‎

Lines changed: 107 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,107 @@
1+
<!-- webkit-test-runner [ enableMetalShaderValidation=true ] -->
2+
<script src="../../../../resources/js-test.js"></script>
3+
<script>
4+
jsTestIsAsync = true;
5+
6+
// `big` and `tiny` are both declared at @group(0) @binding(0). A function-scope declaration is not
7+
// in scope inside its own initializer, so the `big` in `arrayLength(&big)` is the module-scope
8+
// buffer and this entry point uses both of them, which must be rejected. It used to be accepted:
9+
// the validator resolved that `big` to the local being declared and never counted the buffer, while
10+
// the rewriter counted both and emitted a four field argument buffer struct. The layout derived
11+
// from the same shader collapsed the duplicate down to two slots, so the shader read its 24 byte
12+
// struct out of a 16 byte argument buffer and picked up the bound buffer's size where it expected
13+
// a device pointer.
14+
const duplicateBindingCode = `
15+
@group(0) @binding(0) var<storage, read_write> big : array<u32>;
16+
@group(0) @binding(0) var<storage, read_write> tiny : array<u32>;
17+
18+
@compute @workgroup_size(1)
19+
fn main() {
20+
tiny[0] = 1u;
21+
var big : u32 = arrayLength(&big);
22+
tiny[1] = big;
23+
}`;
24+
25+
// Two resources may share a binding as long as no single entry point uses both, so this must keep
26+
// compiling and running. Without it the test could pass by rejecting everything.
27+
const singleBindingCode = `
28+
@group(0) @binding(0) var<storage, read_write> big : array<u32>;
29+
@group(0) @binding(0) var<storage, read_write> tiny : array<u32>;
30+
31+
@compute @workgroup_size(1)
32+
fn main() {
33+
tiny[0] = arrayLength(&tiny);
34+
}`;
35+
36+
let device;
37+
38+
async function compile(code)
39+
{
40+
device.pushErrorScope('validation');
41+
const module = device.createShaderModule({ code });
42+
const messages = (await module.getCompilationInfo()).messages;
43+
const error = await device.popErrorScope();
44+
return { module, messages: [...messages], error };
45+
}
46+
47+
// 0xBEE0 is the u32 the layout writes into the argument buffer where the shader's struct expects
48+
// the low half of a device pointer, so a surviving mismatch stores through address 0xBEE0.
49+
async function dispatch(module)
50+
{
51+
device.pushErrorScope('validation');
52+
const pipeline = device.createComputePipeline({ layout: 'auto', compute: { module, entryPoint: 'main' } });
53+
const buffer = device.createBuffer({ size: 0xBEE0, usage: GPUBufferUsage.STORAGE });
54+
const bindGroup = device.createBindGroup({
55+
layout: pipeline.getBindGroupLayout(0),
56+
entries: [{ binding: 0, resource: { buffer } }],
57+
});
58+
const encoder = device.createCommandEncoder();
59+
const pass = encoder.beginComputePass();
60+
pass.setPipeline(pipeline);
61+
pass.setBindGroup(0, bindGroup);
62+
pass.dispatchWorkgroups(1);
63+
pass.end();
64+
device.queue.submit([encoder.finish()]);
65+
const error = await device.popErrorScope();
66+
await device.queue.onSubmittedWorkDone();
67+
return error;
68+
}
69+
70+
onload = async () => {
71+
try {
72+
const adapter = await navigator.gpu.requestAdapter();
73+
device = await adapter.requestDevice();
74+
75+
const duplicate = await compile(duplicateBindingCode);
76+
const reported = duplicate.messages.some(message => message.type === 'error'
77+
&& message.message.includes('same resource binding: @group(0) @binding(0)'));
78+
if (reported)
79+
testPassed('the duplicate resource binding was reported');
80+
else
81+
testFailed(`the duplicate resource binding was not reported: ${JSON.stringify(duplicate.messages.map(m => m.message))}`);
82+
83+
// Run it anyway. On an implementation that accepted the module this is where the undersized
84+
// argument buffer is read, which Metal shader validation reports as an invalid device load.
85+
try {
86+
await dispatch(duplicate.module);
87+
} catch (e) {
88+
}
89+
testPassed('dispatching the invalid module did not crash the GPU process');
90+
91+
const single = await compile(singleBindingCode);
92+
const singleCompiled = !single.error && !single.messages.some(message => message.type === 'error');
93+
if (!singleCompiled)
94+
testFailed(`an entry point using only one of the two bindings did not compile: ${single.error ? single.error.message : JSON.stringify(single.messages.map(m => m.message))}`);
95+
else {
96+
const error = await dispatch(single.module);
97+
if (error)
98+
testFailed(`an entry point using only one of the two bindings was rejected: ${error.message}`);
99+
else
100+
testPassed('an entry point using only one of the two bindings still runs');
101+
}
102+
} catch (e) {
103+
testFailed(`${e}`);
104+
}
105+
finishJSTest();
106+
};
107+
</script>

‎LayoutTests/sar3/fast/webgpu/regression/shader-module-declaration-scope-expected.txt‎

Lines changed: 10 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,10 @@
1+
PASS a shadowing local does not hide workgroup storage from a vertex entry point
2+
PASS workgroup storage in a vertex entry point is rejected without shadowing
3+
PASS workgroup storage in a compute entry point still compiles
4+
PASS a shadowing local does not drop the dependency on the declaration it shadows
5+
PASS the same module compiles when the dependency is already declared first
6+
PASS a shadowing local does not drop the dependency on a later module-scope variable
7+
PASS successfullyParsed is true
8+
9+
TEST COMPLETE
10+

‎LayoutTests/sar3/fast/webgpu/regression/shader-module-declaration-scope.html‎

Lines changed: 118 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,118 @@
1+
<script src="../../../../resources/js-test.js"></script>
2+
<script>
3+
jsTestIsAsync = true;
4+
5+
// https://www.w3.org/TR/WGSL/#scope: the scope of a function-scope declaration starts just after
6+
// the declaration, so its own initializer resolves names against the enclosing scope. Two of WGSL's
7+
// visitors introduced the name before visiting the declaration, which made a module-scope
8+
// declaration invisible to them whenever a local shadowed it. This test covers the two
9+
// consequences that are not about resource bindings; repro_186926179.html covers that one.
10+
11+
// The `w` in the initializer is the module-scope workgroup variable, so this vertex entry point uses
12+
// workgroup storage and has to be rejected. It used to be accepted, and the emitted Metal declared
13+
// a threadgroup variable and a threadgroup_barrier inside a [[vertex]] function, which the Metal
14+
// compiler rejects.
15+
const workgroupInVertex = `
16+
var<workgroup> w : u32;
17+
@vertex fn main() -> @builtin(position) vec4<f32> {
18+
var w : u32 = w;
19+
return vec4<f32>(f32(w));
20+
}`;
21+
22+
// The same reference without the shadowing local. This was always rejected, so it shows the check
23+
// itself is present rather than the test passing because every module is refused.
24+
const workgroupInVertexUnshadowed = `
25+
var<workgroup> w : u32;
26+
@vertex fn main() -> @builtin(position) vec4<f32> {
27+
let v : u32 = w;
28+
return vec4<f32>(f32(v));
29+
}`;
30+
31+
// A compute entry point may use workgroup storage, shadowing local or not.
32+
const workgroupInCompute = `
33+
var<workgroup> w : u32;
34+
@group(0) @binding(0) var<storage, read_write> out : array<u32>;
35+
@compute @workgroup_size(1) fn main() {
36+
var w : u32 = w;
37+
out[0] = w;
38+
}`;
39+
40+
// `g` in the initializer is the module-scope function, so `f` depends on `g` and `g` has to be
41+
// sorted before it. The dependency edge used to be dropped, `g` stayed after `f`, and this valid
42+
// module was rejected with "unresolved call target 'g'".
43+
const forwardDependency = `
44+
fn f() -> u32 { var g : u32 = g(); return g; }
45+
fn g() -> u32 { return 7u; }
46+
@compute @workgroup_size(1) fn main() { _ = f(); }`;
47+
48+
// The same module with `g` already ahead of `f`, which compiled either way: it isolates the failure
49+
// above to the sorting rather than to name resolution.
50+
const forwardDependencyPresorted = `
51+
fn g() -> u32 { return 7u; }
52+
fn f() -> u32 { var g : u32 = g(); return g; }
53+
@compute @workgroup_size(1) fn main() { _ = f(); }`;
54+
55+
// The same lost edge for a module-scope variable rather than a function. This was rejected with
56+
// "unresolved identifier 'g'".
57+
const forwardVariableDependency = `
58+
fn f() -> u32 { var g : u32 = g; return g; }
59+
var<private> g : u32 = 7u;
60+
@compute @workgroup_size(1) fn main() { _ = f(); }`;
61+
62+
let device;
63+
64+
async function compile(code)
65+
{
66+
device.pushErrorScope('validation');
67+
const module = device.createShaderModule({ code });
68+
const messages = (await module.getCompilationInfo()).messages;
69+
const error = await device.popErrorScope();
70+
const errors = messages.filter(message => message.type === 'error').map(message => message.message);
71+
if (error && !errors.length)
72+
errors.push(error.message);
73+
return errors;
74+
}
75+
76+
async function shouldReject(description, code, expectedSubstring)
77+
{
78+
const errors = await compile(code);
79+
if (errors.some(message => message.includes(expectedSubstring)))
80+
testPassed(description);
81+
else
82+
testFailed(`${description}: expected an error containing "${expectedSubstring}", got ${JSON.stringify(errors)}`);
83+
}
84+
85+
async function shouldCompile(description, code)
86+
{
87+
const errors = await compile(code);
88+
if (!errors.length)
89+
testPassed(description);
90+
else
91+
testFailed(`${description}: ${JSON.stringify(errors)}`);
92+
}
93+
94+
onload = async () => {
95+
try {
96+
const adapter = await navigator.gpu.requestAdapter();
97+
device = await adapter.requestDevice();
98+
99+
const workgroupInVertexError = 'cannot be used by `vertex` pipeline stage';
100+
await shouldReject('a shadowing local does not hide workgroup storage from a vertex entry point',
101+
workgroupInVertex, workgroupInVertexError);
102+
await shouldReject('workgroup storage in a vertex entry point is rejected without shadowing',
103+
workgroupInVertexUnshadowed, workgroupInVertexError);
104+
await shouldCompile('workgroup storage in a compute entry point still compiles',
105+
workgroupInCompute);
106+
107+
await shouldCompile('a shadowing local does not drop the dependency on the declaration it shadows',
108+
forwardDependency);
109+
await shouldCompile('the same module compiles when the dependency is already declared first',
110+
forwardDependencyPresorted);
111+
await shouldCompile('a shadowing local does not drop the dependency on a later module-scope variable',
112+
forwardVariableDependency);
113+
} catch (e) {
114+
testFailed(`${e}`);
115+
}
116+
finishJSTest();
117+
};
118+
</script>

‎Source/WebGPU/WGSL/GlobalSorting.cpp‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -216,14 +216,14 @@ GraphBuilder::GraphBuilder(Graph& graph, Graph::Node& node)
216216

217217
void GraphBuilder::visit(AST::Parameter& parameter)
218218
{
219-
introduceVariable(parameter.name());
220219
Base::visit(parameter);
220+
introduceVariable(parameter.name());
221221
}
222222

223223
void GraphBuilder::visit(AST::VariableStatement& variable)
224224
{
225-
introduceVariable(variable.variable().name());
226225
Base::visit(variable);
226+
introduceVariable(variable.variable().name());
227227
}
228228

229229
void GraphBuilder::visit(AST::Expression& expression)

‎Source/WebGPU/WGSL/GlobalVariableRewriter.cpp‎

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1516,7 +1516,8 @@ auto RewriteGlobalVariables::determineUsedGlobals(const CallGraph::EntryPoint& e
15161516
auto binding = global.resource->binding;
15171517
auto groupResult = usedGlobals.resources.add(group, IndexMap<const Global*>());
15181518
auto bindingResult = groupResult.iterator->value.add(binding, &global);
1519-
ASSERT_UNUSED(bindingResult, bindingResult.isNewEntry);
1519+
if (!bindingResult.isNewEntry) [[unlikely]]
1520+
return makeUnexpected(Error(makeString("entry point '"_s, entryPoint.originalName, "' uses two variables with the same resource binding: @group("_s, group, ") @binding("_s, binding, ')'), function.span()));
15201521
}
15211522

15221523
m_shaderModule.addOverrideValidation([span = function.span(), variables = WTF::move(workgroupVariables), maximumCombinedWorkgroupVariablesSize](auto&) -> std::optional<Error> {

‎Source/WebGPU/WGSL/IOValidator.cpp‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -157,14 +157,14 @@ void IOValidator::visit(AST::Function& function)
157157

158158
void IOValidator::visit(AST::VariableStatement& statement)
159159
{
160-
introduceVariable(statement.variable().name(), nullptr);
161160
AST::Visitor::visit(statement);
161+
introduceVariable(statement.variable().name(), nullptr);
162162
}
163163

164164
void IOValidator::visit(AST::Parameter& parameter)
165165
{
166-
introduceVariable(parameter.name(), nullptr);
167166
AST::Visitor::visit(parameter);
167+
introduceVariable(parameter.name(), nullptr);
168168
}
169169

170170
void IOValidator::visit(AST::IdentifierExpression& identifier)

‎Tools/TestWebKitAPI/Tests/WGSL/TypeCheckingTests.mm‎

Lines changed: 76 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -42,6 +42,8 @@ inline void expectTypeError(const String& wgsl, const String& errorMessage)
4242
.supportedFeatures = { "shader-f16"_s, "clip-distances"_s, "primitive-index"_s }
4343
});
4444
EXPECT_TRUE(std::holds_alternative<WGSL::FailedCheck>(result));
45+
if (!std::holds_alternative<WGSL::FailedCheck>(result))
46+
return;
4547
auto error = std::get<WGSL::FailedCheck>(result);
4648
EXPECT_EQ(error.errors.size(), 1u);
4749
EXPECT_EQ(error.errors[0].message(), errorMessage);
@@ -721,6 +723,80 @@ inline void expectNoError(const String& wgsl)
721723
"uniform address space requires that the number of bytes between S1::y3 and S1::z3 must be at least 16 bytes, but it is 8 bytes"_s);
722724
}
723725

726+
inline void expectStaticCheckErrorContaining(const String& wgsl, const String& errorMessageSubstring)
727+
{
728+
auto result = staticCheck(wgsl);
729+
EXPECT_TRUE(std::holds_alternative<WGSL::FailedCheck>(result));
730+
if (!std::holds_alternative<WGSL::FailedCheck>(result))
731+
return;
732+
auto error = std::get<WGSL::FailedCheck>(result);
733+
EXPECT_EQ(error.errors.size(), 1u);
734+
EXPECT_TRUE(error.errors[0].message().contains(errorMessageSubstring));
735+
}
736+
737+
TEST(WGSLTypeCheckingTests, ResourceBinding)
738+
{
739+
expectStaticCheckErrorContaining(
740+
"@group(0) @binding(0) var<storage, read_write> big: array<u32>;"
741+
"@group(0) @binding(0) var<storage, read_write> tiny: array<u32>;"
742+
"@compute @workgroup_size(1) fn main() {"
743+
"tiny[0] = 1u;"
744+
"var big: u32 = arrayLength(&big);"
745+
"tiny[1] = big;"
746+
"}"_s,
747+
"both which use the same resource binding: @group(0) @binding(0)"_s);
748+
749+
expectStaticCheckErrorContaining(
750+
"@group(0) @binding(0) var<storage, read_write> big: array<u32>;"
751+
"@group(0) @binding(0) var<storage, read_write> tiny: array<u32>;"
752+
"@compute @workgroup_size(1) fn main() {"
753+
"tiny[0] = arrayLength(&big);"
754+
"}"_s,
755+
"both which use the same resource binding: @group(0) @binding(0)"_s);
756+
757+
expectNoError(
758+
"@group(0) @binding(0) var<storage, read_write> big: array<u32>;"
759+
"@group(0) @binding(0) var<storage, read_write> tiny: array<u32>;"
760+
"@compute @workgroup_size(1) fn main() {"
761+
"tiny[0] = arrayLength(&tiny);"
762+
"}"_s);
763+
764+
expectNoError(
765+
"@group(0) @binding(0) var<storage, read_write> big: array<u32>;"
766+
"@group(0) @binding(0) var<storage, read_write> tiny: array<u32>;"
767+
"@compute @workgroup_size(1) fn main() {"
768+
"var big: u32 = 0u;"
769+
"big = big + 1u;"
770+
"tiny[0] = big;"
771+
"}"_s);
772+
}
773+
774+
TEST(WGSLTypeCheckingTests, DeclarationScope)
775+
{
776+
expectTypeError(
777+
"var<workgroup> w: u32;"
778+
"@vertex fn main() -> @builtin(position) vec4<f32> {"
779+
"var w: u32 = w;"
780+
"return vec4<f32>(f32(w));"
781+
"}"_s,
782+
"var with `workgroup` address space cannot be used by `vertex` pipeline stage"_s);
783+
784+
expectNoError(
785+
"fn f() -> u32 { var g: u32 = g(); return g; }"
786+
"fn g() -> u32 { return 7u; }"
787+
"@compute @workgroup_size(1) fn main() { _ = f(); }"_s);
788+
789+
expectNoError(
790+
"fn f() -> u32 { var g: u32 = g; return g; }"
791+
"var<private> g: u32 = 7u;"
792+
"@compute @workgroup_size(1) fn main() { _ = f(); }"_s);
793+
794+
expectTypeError(
795+
"fn f() -> u32 { var f: u32 = f(); return f; }"
796+
"@compute @workgroup_size(1) fn main() { _ = f(); }"_s,
797+
"encountered a dependency cycle: f -> f"_s);
798+
}
799+
724800
TEST(WGSLTypeCheckingTests, Return)
725801
{
726802
expectTypeError("fn noReturn() { return 0i; }"_s, "return statement type does not match its function return type, returned 'i32', expected 'void'"_s);

0 commit comments

Comments
 (0)