[10] WGSL GlobalVariableRewriter iterator-invalidation UAF
Severity: Medium | Component: WebGPU WGSL compiler — GlobalVariableRewriter | efa18a8
Rated Medium because the observable effect is a use-after-free on HashMap backing storage during WGSL compilation — reachable from web content via the WebGPU API — but exploitation requires controlling heap state at rehash time, which is non-trivial given the HashMap's own reallocation behavior.
HashMap modified while iterating in GlobalVariableRewriter. While propagating the list of length parameters from callee to callers, the code mutated the HashMap during iteration. The fix defers insertions into a separate vector.
Source/WebGPU/WGSL/GlobalVariableRewriter.cpp
+ Vector<std::pair<AST::Function*, String>> pendingLengthParameters;
for (auto& lengthParameter : it->value) {
...
for (auto& [caller, call] : callee.callSites) {
...
- auto result = m_lengthParameters.add(caller, ListHashSet<String> { });
- result.iterator->value.add(identifier);
+ pendingLengthParameters.append({ caller, identifier });
...
}
}
+
+ for (auto& [caller, lengthParameter] : pendingLengthParameters) {
+ auto result = m_lengthParameters.add(caller, ListHashSet<String> { });
+ result.iterator->value.add(lengthParameter);
+ }
Tools/TestWebKitAPI/Tests/WGSL/shaders/fuzz-166715941.wgsl
+fn fn0(a0: ptr<storage, array<vec4i>, read_write>) {
+ var a = a0[0];
+}
+fn fn1() { fn0(&buffer32); }
+fn fn2() { fn1(); fn0(&buffer32); }
+fn fn3() { fn2(); fn0(&buffer32); }
+fn fn4(a0: ptr<storage, array<i32>, read>) {
+ var a = buffer32[0];
+ var b = a0[0];
+}
+@compute @workgroup_size(3, 1, 1)
+fn fn5() { fn4(&buffer36); fn3(); }
Patch Details
In RewriteGlobalVariables::visitCallee(), the fix defers mutations to the m_lengthParameters HashMap until after iteration completes. Previously, inside a loop iterating over it->value, the code called m_lengthParameters.add(caller, ...), which could insert new keys and trigger a rehash, invalidating the outer iterator.
HashMap iterator invalidation due to insertion during iteration in a compiler pass.
Background
WTF HashMap is WebKit's hash map implementation. Inserting a new key may trigger a rehash that reallocates the backing store to a larger buffer, freeing the old one and invalidating all iterators. The WGSL GlobalVariableRewriter is a compiler pass that transforms WGSL global variable accesses into explicit function parameters for Metal code generation. It walks the call graph bottom-up, propagating which globals (and their array lengths) each function needs, threading them through call sites as additional arguments. arrayLength() in WGSL is a built-in that returns the runtime size of a runtime-sized array in a storage buffer — the compiler must pass the buffer's byte length as an extra parameter through the call chain, tracked in m_lengthParameters.
Analysis
Before the fix, visitCallee() iterated over m_lengthParameters[callee.target] via an iterator it into the m_lengthParameters HashMap (as strongly implied by the fix pattern and data types). Inside this loop, for each caller's call site, the code called m_lengthParameters.add(caller, ...) to propagate length parameters upward through the call graph. Inserting a new key into a WTF HashMap can trigger a rehash, which reallocates the backing store and invalidates all existing iterators — including it, which is still being used in the outer loop. After rehash, the outer loop dereferences a dangling iterator pointing into the old (freed) backing store.
Aaa Aaaaaa Aaaa Aaaa Aaaaaaaaaaaa Aaa Aaaaaaaa Aaaa Aaaaa Aaaaaa Aaaaa Aaaaa Aaaaa Aaaaaa Aaaa Aaaaaaaaaa Aaa Aa Aaaaaaaaaaaaaaaaaa Aaaaaaa Aaa Aaaaa Aaaaaa Aaaaaa Aaaaaaa Aaaaa a Aaaaa a Aaaaa Aaaaa Aaaaaaaaaaaa Aaaa Aaa Aaaaaaa Aaaa Aaaaaa Aaaaaaa Aaa Aaaaaa Aaaaaaaaa Aaaaaa Aaa Aaaaa Aaaaaaa Aaaaa Aaa Aaaaa Aaaaaaaa Aa Aaaaaaaaaaa Aaa Aaaaaaaaaa Aaaaa Aaaaaa Aaaaa Aaaaaaa
Aaa Aaa Aa Aaaaaaaaa Aaaa Aaa Aaaaaaa Aaa Aaa Aaaaaa Aaa Aa Aaaaaaaaaa a Aaaaaaa Aaaa Aaaaaaa Aaaaaaaaaaaa Aaaaa Aaaaaaa Aaaaaaaaaaa Aaa Aaaa Aaaaa Aa Aaa Aaaa Aa Aaaaaa Aa Aaaa Aaa Aaaaa Aaaaaaa Aaaaa Aaaaaaaa Aaaaaaaaaaaaaaaaaaa Aaaaa Aaaaa Aaa Aaaaaaaa Aaaa Aaaaaaaaaa Aa Aaaaaaaa Aaaaaaaa Aaa Aaaaaaaaa Aaaaaa Aaaa Aa Aaaaaaaaa Aaa Aaaaaaaaaaaa Aa Aaa Aaaaaaaaa Aaa Aaaaaaaaaaaa Aaa Aaaaaaaaaaa Aaaaa Aaa Aaaaa Aaaaaaa
Aaaa Aaaaaaaaaaaaa Aaaaaaa Aaaaaa Aaaaaa Aaaaaa Aaa Aaaaaaa Aaaa Aaaaaaaa Aaaa Aaaaaaaa Aa Aaaaaaaa Aaa Aaaaaa a Aaaa Aaaaaa Aaaa Aaa Aaaaa Aaaa Aaaaa Aaaaaaaaa Aaaaa Aaaaaaa Aaa Aaaaaaa Aaaaaa Aaaaaa Aaaaaaaaaa Aaaaaaa Aaa Aaaaaaaa Aa Aaaa Aaaa Aaaaa Aaaaaaa Aaaaaaaaa Aa Aaaa Aaaaaa Aaaa Aaaaa Aaaaa Aaaaaaaaa Aaaaaaaaaaa Aaaaaaaaaaaa Aaaaaaaaa a Aaaaaaaaa Aaaaaa Aaaa Aaaaaaaa Aaaaaaaaaaaaa Aa a Aaaaaaaaaaaa Aaaaaa Aaaaaaaaaa Aaaaaaaaa Aaaaaa Aaa Aaaaaaaa Aaaaaaaa
Aaaa Aaaaaaaa Aaaaaa Aaaa Aaaaaaaaa Aaaaaaaaaaa Aaaaaaa Aaa Aaaa Aaaaa Aaa Aaaaaaaaaaaa Aaaaa Aa Aaaa Aaaaa Aa Aaaa Aaa Aaaaaaaaaaa Aaaa Aaaaaaaaa Aaaaa Aa Aaaaaa Aaa Aaaa Aaaa Aaaaaaaaa Aa Aa Aaaaaaa Aaaaa
Aaaaaaaaa Aaa Aaaaaaaa Aaaaaaaa Aa Aaaaaaaa Aaaa Aa Aaaaaaaa Aaaa Aaaaaaaaaaaaaaaaaaaa Aaa Aaa Aaaaa Aaaaaa Aaaaaaaa Aa Aaaaaaaaaaaa Aaa Aaaaaaaa Aaaa Aaa Aaa Aaaaaaa Aaa Aaaaaaaa Aaaa Aaa Aaaaaaaaaa Aaa Aaaa Aaaaaaaaaaaaaaaaaaaaa Aaaaaaaaa Aa Aaaaaaaa Aaaaaaaaa Aa Aaa Aaaaaa
🔒Explores the memory corruption mechanism and whether heap conditions could escalate beyond a crash
Subscribe to read more
Audit directions
a Aaaaaaaaa Aa Aaaaaaa Aaaaaaaa Aaaaaa Aaaaaaaaa Aa Aaaaaaaa Aaaaaa Aaaa Aaaaaaaaa Aaaaaaaaaaa Aaaaaaa a Aaaa Aaaaaaaa Aaa Aaaaaaaaa Aaaaa Aaa Aaaaaaa Aaaa Aaaaaaaaaaaaaaaaaaaaaaaaaa Aaa Aaaaaa Aaa Aaaaa Aaaa Aaaaaaaaaaaaaaaaaaaaa Aaaaa Aaa a Aaaaaaaaaaa Aaa Aaaa Aaaaa Aaaaa Aaaa Aaaaaaaa Aaaaaa Aa Aaaaaaaaaaaaaaaaaaaaa Aaaa Aaaa Aaaaaaaaaaaaaaaaaaa Aaaaaaaaaa Aaa Aaaaaa Aaaaaa Aaaaaa Aaaa Aaa Aaaaaaa Aa Aaaaaaa Aaaaa Aaaaaa Aaaaa Aaaa Aaaaaaa Aaaa Aaa Aaaa Aaaaaaaaa Aaaaaaa
a Aaaaaaaaaa Aaaaaaaaaaaa Aa Aaa Aaaaaaaaaa Aaa Aa Aaaaaaaaaaaaaaaaaaa Aaaaaaaaa Aaaaaa Aaaa Aaaaaaaaa Aaaaaaaaaa Aaa Aaaaaaa Aaaa Aaa Aaaaaaa Aaaaaaaa Aaaaaaaaa Aaaaaaaaaa Aaaaaa Aaaaaaaaaaa Aaaaa Aaa Aaaaaa Aaaa Aaaa Aaaaaaaa a Aaaaaaa Aaaaa Aaaaaaaaaaaaa Aaaaaaaaa Aaa Aaaa Aa Aaa Aaaa Aaaaa Aaaaa Aaaa Aaaaaaaaaaaaaaaaaaaaa Aaa Aaaaaaaaaaaaaaaaaaaaaaaaaaaa Aaaaa Aaaaaaaa Aaaaaa Aaaaaaaaaa Aaaaaaaaa Aaaaaaaaaaaa Aaaaaaaa Aaaaaaa Aaaa Aaaaaaa
a Aaaaaaaaaaa Aaaa Aaaaaaaaaaa Aaaaaa Aaaa Aaaaaaa Aaaaa Aaaaaaa Aaaaaaa Aa Aaa Aaaaaaaaa Aaaaaa Aaaaaaaaa Aaaaa Aaaaa Aaa Aaaaaaaa Aa Aaaa Aaaaa Aaaaaa Aaaaaa Aaaaaaaaaa Aaa Aaaaaaaaa Aaaaa Aaaaaa Aaaaa Aaaaaa Aaaaaa Aaaaaa Aa Aaaaaaaaaaaaa Aaaaaaa Aaaaaaaa Aaaaaaaa Aaa Aaaaa Aa Aaaaaaa
🔒Multiple audit patterns identified for similar iterator-invalidation bugs across WebKit compiler passes
Subscribe to read more