← All issues

[4] WGSL Override Validation Crash via Pack Wrapper Misclassification

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

web content에서 도달 가능한 shader 컴파일 경로에서 process crash(DoS)가 매번 재현 가능하게 발생하며, memory corruption이나 권한 상승 경로는 확인되지 않아 Medium으로 평가합니다. crash는 constant evaluator의 type dispatch 로직에서 발생하며, 오염된 메모리를 사용하는 과정에서 일어나는 것이 아닙니다.

GlobalVariableRewriter::pack()이 생성하는 __pack/__unpack wrapper CallExpressionresolvedTarget을 설정하고, evaluate()의 struct constructor 경로를 resolvedTarget null 검사로 보호합니다.

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

두 가지 변경이 crash를 수정합니다. 먼저 GlobalVariableRewriter.cpp에서는 RewriteGlobalVariables::pack()이 생성하는 __pack/__unpack wrapper CallExpressioncall.m_resolvedTarget = operation을 설정하도록 변경되었습니다. WGSL.cpp에서는 evaluate() 함수의 dispatch 구조가 재편되었습니다. inferredType()을 기반으로 ConstantStruct를 생성하는 struct constructor 경로가 이제 if (!function) 조건 안으로 이동하여, resolvedTarget이 null인 경우에만 실행됩니다. 패치 이전에는 struct inferredType()을 가진 CallExpression이라면 __pack wrapper를 포함하여 resolvedTarget이 null이 아닌 경우에도 struct constructor 경로로 진입하는 상황이 가능했습니다.

WGSL은 WebGPU API에서 사용하는 shading language입니다. web content는 device.createShaderModule()을 통해 shader 코드를 GPU 실행용으로 제출합니다. WGSL의 override 선언은 pipeline 생성 시 evaluate()를 통해 값이 결정되는 pipeline-overridable constant입니다. GlobalVariableRewriter 패스는 packing이 필요한 storage buffer 접근(예: vec3u__packed_vec3<u32>로 pack하는 경우)을 __pack(...)/__unpack(...) CallExpression 노드로 감싸는 방식으로 rewrite합니다. CallExpressioninferredType()(결과 타입)과 resolvedTarget()(호출 대상 함수) 두 가지 정보를 모두 갖습니다. WGSL에서 struct constructor는 resolvedTarget이 null이고 struct inferredType을 갖는 CallExpression으로 표현되며, 이것이 evaluate()가 dispatch를 결정할 때 사용하는 구분 기준입니다.

근본 원인은 reference capture 위험과 inference 기반 dispatch의 결합입니다. GlobalVariableRewriter::pack()&expression을 reference로 캡처하는 override validation lambda를 등록합니다. 이후 visitAndReplace가 해당 expression을 __pack(...) 노드로 감싸면, lambda가 보유하던 reference는 원래 CallExpression이 아닌 wrapper를 가리키게 됩니다. __pack wrapper는 pack()resolvedTarget()을 설정하지 않았기 때문에 null resolvedTarget()을 가지며, 내부 argument의 타입을 그대로 상속하므로 원래와 동일한 struct inferredType()을 갖습니다.

override validation 과정에서 이 wrapper에 대해 evaluate()가 호출됩니다. resolvedTarget()이 null이고 inferredType()이 struct 타입을 반환하기 때문에, evaluate()는 struct constructor 경로로 진입하게 됩니다. 이 경로에서 구조가 맞지 않는 ConstantStruct가 생성되고, 결국 convertValueImpl에서 crash가 발생합니다. 테스트 케이스는 bufstorage buffer이고 ov가 override인 buf = S(vec3u(u32(ov))) 형태의 shader입니다. S(...) struct constructor가 packed storage buffer를 대상으로 하기 때문에 __pack(...)으로 감싸지고, override expression ov가 감싸진 expression에 대한 constant evaluation을 강제합니다.

🔒

상세 취약점 분석, 공격 가능성 평가, 보안 영향 분석이 포함되어 있습니다

더 확인하려면 구독해 주세요

🔒

이 취약점 패턴의 변종을 찾기 위한 구체적인 탐색 방향이 포함되어 있습니다

더 확인하려면 구독해 주세요