← All issues

[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(); }

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.

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.

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.

🔒

Explores the memory corruption mechanism and whether heap conditions could escalate beyond a crash

Subscribe to read more

🔒

Multiple audit patterns identified for similar iterator-invalidation bugs across WebKit compiler passes

Subscribe to read more