← All issues

[4] WGSL Override Validation Crash via Pack Wrapper Misclassification

Severity: Medium | Component: WebGPU WGSL Compiler (GlobalVariableRewriter, constant evaluator) | 96c9f5e

Rated Medium because the observable effect is a reliable process crash (DoS) in shader compilation reachable from web content, with no memory corruption or escalation path identified — the crash occurs in the constant evaluator's type-dispatch logic, not in a use of corrupted memory.

Set resolvedTarget on __pack/__unpack wrapper CallExpressions created by GlobalVariableRewriter::pack(), and guard the struct constructor path in evaluate() behind a null resolvedTarget check.

Source/WebGPU/WGSL/GlobalVariableRewriter.cpp

call.m_inferredType = argument.inferredType();
+ call.m_resolvedTarget = operation;
m_shaderModule.replace(expression, call);

Source/WebGPU/WGSL/WGSL.cpp

- if (auto* structType = std::get_if<Types::Struct>(expression.inferredType())) {
- HashMap<String, ConstantValue> constantFields;
- for (unsigned i = 0; i < argumentCount; ++i) {
- auto& argument = arguments[i];
- auto& member = structType->structure.members()[i];
- constantFields.set(member.originalName(), argument);
+ if (!function) {
+ if (auto* structType = std::get_if<Types::Struct>(expression.inferredType())) {
+ HashMap<String, ConstantValue> constantFields;
+ for (unsigned i = 0; i < argumentCount; ++i) {
+ auto& argument = arguments[i];
+ auto& member = structType->structure.members()[i];
+ constantFields.set(member.originalName(), argument);
+ }
+ return ConstantStruct { WTF::move(constantFields) };
}
- return ConstantStruct { WTF::move(constantFields) };
- }
-
- if (!function)
return std::nullopt;
+ }

LayoutTests/fast/webgpu/nocrash/fuzz-172084012.html

+struct S { x: vec3u }
+override ov: f32 = 1.0;
+@group(0) @binding(0) var<storage, read_write> buf: S;
+@compute @workgroup_size(1)
+fn main() { buf = S(vec3u(u32(ov))); }

Two changes fix the crash. In GlobalVariableRewriter.cpp, RewriteGlobalVariables::pack() now sets call.m_resolvedTarget = operation on the __pack/__unpack wrapper CallExpression it creates. In WGSL.cpp, the evaluate() function restructures its dispatch: the struct-constructor path (which builds a ConstantStruct from the expression's inferredType()) is now gated behind if (!function), meaning it only fires when resolvedTarget is null. Previously, any CallExpression with a struct inferredType() — including __pack wrappers — could enter the struct-constructor path even though it had a non-null resolvedTarget.

WGSL is the shading language for the WebGPU API. Web content submits shader code for GPU execution via device.createShaderModule(). WGSL override declarations are pipeline-overridable constants whose values are resolved at pipeline creation time via evaluate(). The GlobalVariableRewriter pass rewrites accesses to storage buffers that require packing (e.g., vec3u packed to __packed_vec3<u32>) by wrapping expressions in __pack(...) / __unpack(...) CallExpression nodes. A CallExpression has both an inferredType() (the result type) and a resolvedTarget() (the function being called). Struct constructors in WGSL are represented as CallExpression nodes with a null resolvedTarget and a struct inferredType — this is the distinguishing characteristic that evaluate() uses for dispatch.

The root cause is a reference-capture hazard combined with inference-based dispatch. GlobalVariableRewriter::pack() registers an override validation lambda that captures &expression by reference. When visitAndReplace subsequently wraps the expression in __pack(...), the lambda's reference now points to the wrapper instead of the original CallExpression. The __pack wrapper has null resolvedTarget() (because pack() did not set it) and the same struct inferredType() as the original (because it inherits the type from the inner argument).

During override validation, evaluate() is called on this wrapper. Because resolvedTarget() is null and inferredType() returns a struct type, evaluate() enters the struct-constructor path. This creates a nested ConstantStruct with mismatched structure, which crashes in convertValueImpl. The test case — a shader with buf = S(vec3u(u32(ov))) where buf is a storage buffer and ov is an override — triggers exactly this path: the struct constructor S(...) is wrapped in __pack(...) because it targets a packed storage buffer, and the override expression ov forces constant evaluation of the wrapped expression.

🔒

Detailed vulnerability analysis & security impact assessment

Subscribe to read more

🔒

Pattern-based audit directions for variant discovery

Subscribe to read more