[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

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

Source/WebGPU/WGSL/tests/valid/concretization.wgsl

@@ -47,3 +47,11 @@ fn testInitializerConcretization() {
   // CHECK-L: vec<unsigned, 2>(0, 0)
   let x2 : vec2<u32> = vec2(0, 0);
 }
+
+@group(0) @binding(0) var<storage, read_write> a: array<atomic<i32>>;
+@compute @workgroup_size(1)
+fn testArrayAccessMaterialization()
+{
+    let i = 0;
+    let x = atomicLoad(&a[i]);
+}