From f425f41fe590130e789e92b8ad71e0d54d69a2a6 Mon Sep 17 00:00:00 2001 From: Tadeu Zagallo Date: Fri, 20 Oct 2023 12:59:12 -0700 Subject: [PATCH] [WGSL] Concretizing atomic types shouldn't assert https://bugs.webkit.org/show_bug.cgi?id=263402 rdar://117223889 Reviewed by Dan Glastonbury. When concretization was originally implemented it was only used when initializing variables, and since variables can't initialized with atomic values, it just asserted that we never tried to concretize an atomic type. However, we also need to concretize array elements when accessing an array with a non-constant index. That was implemented in 269416@main, and now we can crash if indexing into array of atomics. The fix is simple, concretizing an atomic type should just return the same type. * Source/WebGPU/WGSL/Constraints.cpp: (WGSL::concretize): * Source/WebGPU/WGSL/tests/valid/concretization.wgsl: Canonical link: https://commits.webkit.org/269579@main --- Source/WebGPU/WGSL/Constraints.cpp | 6 +++--- Source/WebGPU/WGSL/tests/valid/concretization.wgsl | 8 ++++++++ 2 files changed, 11 insertions(+), 3 deletions(-) diff --git a/Source/WebGPU/WGSL/Constraints.cpp b/Source/WebGPU/WGSL/Constraints.cpp index 4f0e01b44807..4d6868d773a4 100644 --- a/Source/WebGPU/WGSL/Constraints.cpp +++ b/Source/WebGPU/WGSL/Constraints.cpp @@ -170,6 +170,9 @@ const Type* concretize(const Type* type, TypeStore& types) [&](const Bottom&) -> const Type* { return type; }, + [&](const Atomic&) -> const Type* { + return type; + }, [&](const Function&) -> const Type* { RELEASE_ASSERT_NOT_REACHED(); }, @@ -185,9 +188,6 @@ const Type* concretize(const Type* type, TypeStore& types) [&](const Reference&) -> const Type* { RELEASE_ASSERT_NOT_REACHED(); }, - [&](const Atomic&) -> const Type* { - RELEASE_ASSERT_NOT_REACHED(); - }, [&](const TypeConstructor&) -> const Type* { RELEASE_ASSERT_NOT_REACHED(); }); diff --git a/Source/WebGPU/WGSL/tests/valid/concretization.wgsl b/Source/WebGPU/WGSL/tests/valid/concretization.wgsl index ecca0e89f976..c48bc2c790f4 100644 --- a/Source/WebGPU/WGSL/tests/valid/concretization.wgsl +++ b/Source/WebGPU/WGSL/tests/valid/concretization.wgsl @@ -47,3 +47,11 @@ fn testInitializerConcretization() { // CHECK-L: vec(0, 0) let x2 : vec2 = vec2(0, 0); } + +@group(0) @binding(0) var a: array>; +@compute @workgroup_size(1) +fn testArrayAccessMaterialization() +{ + let i = 0; + let x = atomicLoad(&a[i]); +}