[naga] Ensure wgsl snapshot code is reachable from an entry point (#7674)

naga/tests/in/wgsl/6220-break-from-loop.wgsl

@@ -41,3 +41,8 @@ fn break_from_loop() {
       break;
     }
 }
+
+@compute @workgroup_size(1)
+fn main() {
+    break_from_loop();
+}

naga/tests/in/wgsl/7048-multiple-dynamic-1.wgsl

@@ -1,3 +1,4 @@
+@compute @workgroup_size(1)
 fn f() {
     let b = array<vec3f, 2>();
     var poly = vec4f(0);

naga/tests/in/wgsl/7048-multiple-dynamic-3.wgsl

@@ -14,3 +14,9 @@ fn foobar(normals: array<vec3f, 12>, count: u32) -> QEFResult {
 
         return QEFResult(0.0, vec3(0.0));
 }
+
+@fragment
+fn main() {
+    var arr: array<vec3f, 12>;
+    foobar(arr, 1);
+}

naga/tests/in/wgsl/abstract-types-atomic.wgsl

@@ -32,3 +32,9 @@ fn test_atomic_u32() {
   _ = atomicMax(&atomic_u32, 1);
   _ = atomicExchange(&atomic_u32, 1);
 }
+
+@compute @workgroup_size(1)
+fn main() {
+    test_atomic_i32();
+    test_atomic_u32();
+}

naga/tests/in/wgsl/abstract-types-builtins.wgsl

@@ -1,3 +1,4 @@
+@compute @workgroup_size(1)
 fn f() {
   // For calls that return abstract types, we can't write the type we
   // actually expect, but we can at least require an automatic

naga/tests/in/wgsl/abstract-types-function-calls.wgsl

@@ -22,6 +22,7 @@ const const_mat_af = mat2x2(vec2(0.0), vec2(0.0));
 const const_arr_af = array(0.0, 0.0);
 const const_arr_ai = array(0, 0);
 
+@compute @workgroup_size(1)
 fn main() {
     func_f(0.0);
     func_f(0);

naga/tests/in/wgsl/abstract-types-let.wgsl

@@ -113,3 +113,9 @@ fn mixed_constant_and_runtime_arguments() {
     let xvusu: vec2<u32> = vec2(u);
     let xvfsf: vec2<f32> = vec2(f);
 }
+
+@compute @workgroup_size(1)
+fn main() {
+    all_constant_arguments();
+    mixed_constant_and_runtime_arguments();
+}

naga/tests/in/wgsl/abstract-types-operators.wgsl

@@ -84,3 +84,10 @@ fn wgpu_4435() {
     let x = 1;
     let y = a[x-1];
 }
+
+@compute @workgroup_size(1)
+fn main() {
+    runtime_values();
+    wgpu_4445();
+    wgpu_4435();
+}

naga/tests/in/wgsl/abstract-types-return.wgsl

@@ -1,6 +1,3 @@
-@compute @workgroup_size(1)
-fn main() {}
-
 fn return_i32_ai() -> i32 {
     return 1;
 }
@@ -34,3 +31,15 @@ fn return_vec2f32_const_ai() -> vec2<f32> {
     const vec_one = vec2(1);
     return vec_one;
 }
+
+@compute @workgroup_size(1)
+fn main() {
+    return_i32_ai();
+    return_u32_ai();
+    return_f32_ai();
+    return_f32_af();
+    return_vec2f32_ai();
+    return_arrf32_ai();
+    return_const_f32_const_ai();
+    return_vec2f32_const_ai();
+}

naga/tests/in/wgsl/abstract-types-texture.wgsl

@@ -24,3 +24,10 @@ fn depth() {
 fn storage() {
   textureStore(st, vec2(0,1), vec4(2,3,4,5));
 }
+
+@fragment
+fn main() {
+    color();
+    depth();
+    storage();
+}

naga/tests/in/wgsl/abstract-types-var.wgsl

@@ -275,3 +275,9 @@ fn mixed_constant_and_runtime_arguments() {
     xvusu = vec2(u);
     xvfsf = vec2(f);
 }
+
+@compute @workgroup_size(1)
+fn main() {
+    all_constant_arguments();
+    mixed_constant_and_runtime_arguments();
+}

naga/tests/in/wgsl/access.wgsl

@@ -159,7 +159,6 @@ fn assign_array_through_ptr_fn(foo: ptr<function, array<vec4<f32>, 2>>) {
     *foo = array<vec4<f32>, 2>(vec4(1.0), vec4(2.0));
 }
 
-@compute @workgroup_size(1)
 fn assign_through_ptr() {
     var val = 33u;
     assign_through_ptr_fn(&val);
@@ -188,7 +187,6 @@ fn assign_to_arg_ptr_array_element(p: ptr<function, array<u32, 4>>) {
   (*p)[1] = 10u;
 }
 
-@compute @workgroup_size(1)
 fn assign_to_ptr_components() {
    var s1: AssignToMember;
    assign_to_arg_ptr_member(&s1);
@@ -242,3 +240,13 @@ fn var_members_of_members() -> i32 {
 
     return thing.om_nom_nom.delicious;
 }
+
+@compute @workgroup_size(1)
+fn foo_compute() {
+    assign_through_ptr();
+    assign_to_ptr_components();
+    index_ptr(true);
+    member_ptr();
+    let_members_of_members();
+    var_members_of_members();
+}

naga/tests/in/wgsl/bounds-check-image-restrict-depth.toml

@@ -0,0 +1,14 @@
+targets = "SPIRV | METAL"
+
+[bounds_check_policies]
+image_load = "Restrict"
+
+[msl]
+fake_missing_bindings = true
+lang_version = [1, 2]
+spirv_cross_compatibility = false
+zero_initialize_workgroup_memory = true
+
+[spv]
+debug = true
+version = [1, 1]

naga/tests/in/wgsl/bounds-check-image-restrict-depth.wgsl

@@ -0,0 +1,36 @@
+// Cases from bounds-check-image-restrict that GLSL does not yet support.
+
+@group(0) @binding(0)
+var image_depth_2d: texture_depth_2d;
+
+fn test_textureLoad_depth_2d(coords: vec2<i32>, level: i32) -> f32 {
+   return textureLoad(image_depth_2d, coords, level);
+}
+
+@group(0) @binding(1)
+var image_depth_2d_array: texture_depth_2d_array;
+
+fn test_textureLoad_depth_2d_array_u(coords: vec2<i32>, index: u32, level: i32) -> f32 {
+   return textureLoad(image_depth_2d_array, coords, index, level);
+}
+
+fn test_textureLoad_depth_2d_array_s(coords: vec2<i32>, index: i32, level: i32) -> f32 {
+   return textureLoad(image_depth_2d_array, coords, index, level);
+}
+
+@group(0) @binding(2)
+var image_depth_multisampled_2d: texture_depth_multisampled_2d;
+
+fn test_textureLoad_depth_multisampled_2d(coords: vec2<i32>, _sample: i32) -> f32 {
+   return textureLoad(image_depth_multisampled_2d, coords, _sample);
+}
+
+@fragment
+fn fragment_shader() -> @location(0) vec4<f32> {
+    test_textureLoad_depth_2d(vec2<i32>(), 0);
+    test_textureLoad_depth_2d_array_u(vec2<i32>(), 0u, 0);
+    test_textureLoad_depth_2d_array_s(vec2<i32>(), 0, 0);
+    test_textureLoad_depth_multisampled_2d(vec2<i32>(), 0);
+
+    return vec4<f32>(0.,0.,0.,0.);
+}

naga/tests/in/wgsl/bounds-check-image-restrict.wgsl

@@ -38,45 +38,20 @@ fn test_textureLoad_multisampled_2d(coords: vec2<i32>, _sample: i32) -> vec4<f32
 }
 
 @group(0) @binding(5)
-var image_depth_2d: texture_depth_2d;
-
-fn test_textureLoad_depth_2d(coords: vec2<i32>, level: i32) -> f32 {
-   return textureLoad(image_depth_2d, coords, level);
-}
-
-@group(0) @binding(6)
-var image_depth_2d_array: texture_depth_2d_array;
-
-fn test_textureLoad_depth_2d_array_u(coords: vec2<i32>, index: u32, level: i32) -> f32 {
-   return textureLoad(image_depth_2d_array, coords, index, level);
-}
-
-fn test_textureLoad_depth_2d_array_s(coords: vec2<i32>, index: i32, level: i32) -> f32 {
-   return textureLoad(image_depth_2d_array, coords, index, level);
-}
-
-@group(0) @binding(7)
-var image_depth_multisampled_2d: texture_depth_multisampled_2d;
-
-fn test_textureLoad_depth_multisampled_2d(coords: vec2<i32>, _sample: i32) -> f32 {
-   return textureLoad(image_depth_multisampled_2d, coords, _sample);
-}
-
-@group(0) @binding(8)
 var image_storage_1d: texture_storage_1d<rgba8unorm, write>;
 
 fn test_textureStore_1d(coords: i32, value: vec4<f32>) {
     textureStore(image_storage_1d, coords, value);
 }
 
-@group(0) @binding(9)
+@group(0) @binding(6)
 var image_storage_2d: texture_storage_2d<rgba8unorm, write>;
 
 fn test_textureStore_2d(coords: vec2<i32>, value: vec4<f32>) {
     textureStore(image_storage_2d, coords, value);
 }
 
-@group(0) @binding(10)
+@group(0) @binding(7)
 var image_storage_2d_array: texture_storage_2d_array<rgba8unorm, write>;
 
 fn test_textureStore_2d_array_u(coords: vec2<i32>, array_index: u32, value: vec4<f32>) {
@@ -87,7 +62,7 @@ fn test_textureStore_2d_array_s(coords: vec2<i32>, array_index: i32, value: vec4
  textureStore(image_storage_2d_array, coords, array_index, value);
 }
 
-@group(0) @binding(11)
+@group(0) @binding(8)
 var image_storage_3d: texture_storage_3d<rgba8unorm, write>;
 
 fn test_textureStore_3d(coords: vec3<i32>, value: vec4<f32>) {
@@ -104,11 +79,6 @@ fn fragment_shader() -> @location(0) vec4<f32> {
     test_textureLoad_2d_array_s(vec2<i32>(), 0, 0);
     test_textureLoad_3d(vec3<i32>(), 0);
     test_textureLoad_multisampled_2d(vec2<i32>(), 0);
-    // Not yet implemented for GLSL:
-    // test_textureLoad_depth_2d(vec2<i32>(), 0);
-    // test_textureLoad_depth_2d_array_u(vec2<i32>(), 0u, 0);
-    // test_textureLoad_depth_2d_array_s(vec2<i32>(), 0, 0);
-    // test_textureLoad_depth_multisampled_2d(vec2<i32>(), 0);
     test_textureStore_1d(0, vec4<f32>());
     test_textureStore_2d(vec2<i32>(), vec4<f32>());
     test_textureStore_2d_array_u(vec2<i32>(), 0u, vec4<f32>());

naga/tests/in/wgsl/bounds-check-image-rzsw-depth.toml

@@ -0,0 +1,18 @@
+targets = "SPIRV | METAL"
+
+[bounds_check_policies]
+image_load = "ReadZeroSkipWrite"
+
+[msl]
+fake_missing_bindings = true
+lang_version = [1, 2]
+spirv_cross_compatibility = false
+zero_initialize_workgroup_memory = true
+
+[glsl]
+version.Desktop = 430
+zero_initialize_workgroup_memory = true
+
+[spv]
+debug = true
+version = [1, 1]

naga/tests/in/wgsl/bounds-check-image-rzsw-depth.wgsl

@@ -0,0 +1,36 @@
+// Cases from bounds-check-image-restrict that GLSL does not yet support.
+
+@group(0) @binding(0)
+var image_depth_2d: texture_depth_2d;
+
+fn test_textureLoad_depth_2d(coords: vec2<i32>, level: i32) -> f32 {
+   return textureLoad(image_depth_2d, coords, level);
+}
+
+@group(0) @binding(1)
+var image_depth_2d_array: texture_depth_2d_array;
+
+fn test_textureLoad_depth_2d_array_u(coords: vec2<i32>, index: u32, level: i32) -> f32 {
+   return textureLoad(image_depth_2d_array, coords, index, level);
+}
+
+fn test_textureLoad_depth_2d_array_s(coords: vec2<i32>, index: i32, level: i32) -> f32 {
+   return textureLoad(image_depth_2d_array, coords, index, level);
+}
+
+@group(0) @binding(2)
+var image_depth_multisampled_2d: texture_depth_multisampled_2d;
+
+fn test_textureLoad_depth_multisampled_2d(coords: vec2<i32>, _sample: i32) -> f32 {
+   return textureLoad(image_depth_multisampled_2d, coords, _sample);
+}
+
+@fragment
+fn fragment_shader() -> @location(0) vec4<f32> {
+    test_textureLoad_depth_2d(vec2<i32>(), 0);
+    test_textureLoad_depth_2d_array_u(vec2<i32>(), 0u, 0);
+    test_textureLoad_depth_2d_array_s(vec2<i32>(), 0, 0);
+    test_textureLoad_depth_multisampled_2d(vec2<i32>(), 0);
+
+    return vec4<f32>(0.,0.,0.,0.);
+}

naga/tests/in/wgsl/bounds-check-image-rzsw.wgsl

@@ -38,45 +38,20 @@ fn test_textureLoad_multisampled_2d(coords: vec2<i32>, _sample: i32) -> vec4<f32
 }
 
 @group(0) @binding(5)
-var image_depth_2d: texture_depth_2d;
-
-fn test_textureLoad_depth_2d(coords: vec2<i32>, level: i32) -> f32 {
-   return textureLoad(image_depth_2d, coords, level);
-}
-
-@group(0) @binding(6)
-var image_depth_2d_array: texture_depth_2d_array;
-
-fn test_textureLoad_depth_2d_array_u(coords: vec2<i32>, index: u32, level: i32) -> f32 {
-   return textureLoad(image_depth_2d_array, coords, index, level);
-}
-
-fn test_textureLoad_depth_2d_array_s(coords: vec2<i32>, index: i32, level: i32) -> f32 {
-   return textureLoad(image_depth_2d_array, coords, index, level);
-}
-
-@group(0) @binding(7)
-var image_depth_multisampled_2d: texture_depth_multisampled_2d;
-
-fn test_textureLoad_depth_multisampled_2d(coords: vec2<i32>, _sample: i32) -> f32 {
-   return textureLoad(image_depth_multisampled_2d, coords, _sample);
-}
-
-@group(0) @binding(8)
 var image_storage_1d: texture_storage_1d<rgba8unorm, write>;
 
 fn test_textureStore_1d(coords: i32, value: vec4<f32>) {
     textureStore(image_storage_1d, coords, value);
 }
 
-@group(0) @binding(9)
+@group(0) @binding(6)
 var image_storage_2d: texture_storage_2d<rgba8unorm, write>;
 
 fn test_textureStore_2d(coords: vec2<i32>, value: vec4<f32>) {
     textureStore(image_storage_2d, coords, value);
 }
 
-@group(0) @binding(10)
+@group(0) @binding(7)
 var image_storage_2d_array: texture_storage_2d_array<rgba8unorm, write>;
 
 fn test_textureStore_2d_array_u(coords: vec2<i32>, array_index: u32, value: vec4<f32>) {
@@ -87,7 +62,7 @@ fn test_textureStore_2d_array_s(coords: vec2<i32>, array_index: i32, value: vec4
  textureStore(image_storage_2d_array, coords, array_index, value);
 }
 
-@group(0) @binding(11)
+@group(0) @binding(8)
 var image_storage_3d: texture_storage_3d<rgba8unorm, write>;
 
 fn test_textureStore_3d(coords: vec3<i32>, value: vec4<f32>) {
@@ -104,11 +79,6 @@ fn fragment_shader() -> @location(0) vec4<f32> {
     test_textureLoad_2d_array_s(vec2<i32>(), 0, 0);
     test_textureLoad_3d(vec3<i32>(), 0);
     test_textureLoad_multisampled_2d(vec2<i32>(), 0);
-    // Not yet implemented for GLSL:
-    // test_textureLoad_depth_2d(vec2<i32>(), 0);
-    // test_textureLoad_depth_2d_array_u(vec2<i32>(), 0u, 0);
-    // test_textureLoad_depth_2d_array_s(vec2<i32>(), 0, 0);
-    // test_textureLoad_depth_multisampled_2d(vec2<i32>(), 0);
     test_textureStore_1d(0, vec4<f32>());
     test_textureStore_2d(vec2<i32>(), vec4<f32>());
     test_textureStore_2d_array_u(vec2<i32>(), 0u, vec4<f32>());

naga/tests/in/wgsl/bounds-check-restrict.wgsl

@@ -78,3 +78,24 @@ fn index_dynamic_array_constant_index() -> f32 {
 fn set_dynamic_array_constant_index(v: f32) {
    globals.d[1000] = v;
 }
+
+@compute @workgroup_size(1)
+fn main() {
+    index_array(1);
+    index_dynamic_array(1);
+    index_vector(1);
+    index_vector_by_value(vec4f(2, 3, 4, 5), 6);
+    index_matrix(1);
+    index_twice(1, 2);
+    index_expensive(1);
+    index_in_bounds();
+    set_array(1, 2.);
+    set_dynamic_array(1, 2.);
+    set_vector(1, 2.);
+    set_matrix(1, vec4f(2, 3, 4, 5));
+    set_index_twice(1, 2, 1.);
+    set_expensive(1, 1.);
+    set_in_bounds(1.);
+    index_dynamic_array_constant_index();
+    set_dynamic_array_constant_index(1.);
+}

naga/tests/in/wgsl/bounds-check-zero-atomic.wgsl

@@ -44,3 +44,14 @@ fn exchange_atomic_dynamic_sized_array_static_index() -> u32 {
    return atomicExchange(&globals.c[1000], 1u);
 }
 
+@compute @workgroup_size(1)
+fn main() {
+    fetch_add_atomic();
+    fetch_add_atomic_static_sized_array(1);
+    fetch_add_atomic_dynamic_sized_array(1);
+    exchange_atomic();
+    exchange_atomic_static_sized_array(1);
+    exchange_atomic_dynamic_sized_array(1);
+    fetch_add_atomic_dynamic_sized_array_static_index();
+    exchange_atomic_dynamic_sized_array_static_index();
+}

naga/tests/in/wgsl/bounds-check-zero.wgsl

@@ -78,3 +78,24 @@ fn index_dynamic_array_constant_index() -> f32 {
 fn set_dynamic_array_constant_index(v: f32) {
    globals.d[1000] = v;
 }
+
+@compute @workgroup_size(1)
+fn main() {
+    index_array(1);
+    index_dynamic_array(1);
+    index_vector(1);
+    index_vector_by_value(vec4f(2, 3, 4, 5), 6);
+    index_matrix(1);
+    index_twice(1, 2);
+    index_expensive(1);
+    index_in_bounds();
+    set_array(1, 2.);
+    set_dynamic_array(1, 2.);
+    set_vector(1, 2.);
+    set_matrix(1, vec4f(2, 3, 4, 5));
+    set_index_twice(1, 2, 1.);
+    set_expensive(1, 1.);
+    set_in_bounds(1.);
+    index_dynamic_array_constant_index();
+    set_dynamic_array_constant_index(1.);
+}

naga/tests/in/wgsl/break-if.wgsl

@@ -1,6 +1,3 @@
-@compute @workgroup_size(1)
-fn main() {}
-
 fn breakIfEmpty() {
     loop {
         continuing {
@@ -42,3 +39,11 @@ fn breakIfSeparateVariable() {
         }
     }
 }
+
+@compute @workgroup_size(1)
+fn main() {
+    breakIfEmpty();
+    breakIfEmptyBody(false);
+    breakIf(false);
+    breakIfSeparateVariable();
+}

naga/tests/in/wgsl/const-exprs.wgsl

@@ -11,8 +11,14 @@ fn main() {
     non_constant_initializers();
     splat_of_constant();
     compose_of_constant();
+    map_texture_kind(1);
     compose_of_splat();
     test_local_const();
+    compose_vector_zero_val_binop();
+    relational();
+    packed_dot_product();
+    test_local_const();
+    abstract_access(1);
 }
 
 // Swizzle the value of nested Compose expressions.

naga/tests/in/wgsl/const_assert.wgsl

@@ -9,6 +9,7 @@ const_assert x == 1i;
 const_assert x > 0u;
 const_assert x < 2.0f;
 
+@compute @workgroup_size(1)
 fn foo() {
   const z = x + y - 2;
   const_assert z > 0; // valid in functions.

naga/tests/in/wgsl/control-flow.wgsl

@@ -1,5 +1,4 @@
-@compute @workgroup_size(1)
+fn control_flow() {
-fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {
     //TODO: execution-only barrier?
     storageBarrier();
     workgroupBarrier();
@@ -217,3 +216,15 @@ fn loop_switch_omit_continue_variable_checks(x: i32, y: i32, z: i32, w: i32) {
         // check needs to be generated here
     }
 }
+
+@compute @workgroup_size(1)
+fn main() {
+    control_flow();
+    switch_default_break(1);
+    switch_case_break();
+    switch_selector_type_conversion();
+    switch_const_expr_case_selectors();
+    loop_switch_continue(1);
+    loop_switch_continue_nesting(1, 2, 3);
+    loop_switch_omit_continue_variable_checks(1, 2, 3, 4);
+}

naga/tests/in/wgsl/conversion-float-to-int-no-f64.wgsl

@@ -107,3 +107,24 @@ fn test_f32_to_i64_vec(f: vec2<f32>) -> vec2<i64> {
 fn test_f32_to_u64_vec(f: vec2<f32>) -> vec2<u64> {
   return vec2<u64>(f);
 }
+
+@compute @workgroup_size(1)
+fn main() {
+    test_const_eval();
+    test_f16_to_i32(1.);
+    test_f16_to_u32(1.);
+    test_f16_to_i64(1.);
+    test_f16_to_u64(1.);
+    test_f32_to_i32(1.);
+    test_f32_to_u32(1.);
+    test_f32_to_i64(1.);
+    test_f32_to_u64(1.);
+    test_f16_to_i32_vec(vec2(1, 2));
+    test_f16_to_u32_vec(vec2(1, 2));
+    test_f16_to_i64_vec(vec2(1, 2));
+    test_f16_to_u64_vec(vec2(1, 2));
+    test_f32_to_i32_vec(vec2(1, 2));
+    test_f32_to_u32_vec(vec2(1, 2));
+    test_f32_to_i64_vec(vec2(1, 2));
+    test_f32_to_u64_vec(vec2(1, 2));
+}

naga/tests/in/wgsl/conversion-float-to-int.wgsl

@@ -144,4 +144,30 @@ fn test_f64_to_u64_vec(f: vec2<f64>) -> vec2<u64> {
 }
 
 @compute @workgroup_size(1)
-fn main() {}
+fn main() {
+    test_const_eval();
+    test_f16_to_i32(1.);
+    test_f16_to_u32(1.);
+    test_f16_to_i64(1.);
+    test_f16_to_u64(1.);
+    test_f32_to_i32(1.);
+    test_f32_to_u32(1.);
+    test_f32_to_i64(1.);
+    test_f32_to_u64(1.);
+    test_f64_to_i32(1.);
+    test_f64_to_u32(1.);
+    test_f64_to_i64(1.);
+    test_f64_to_u64(1.);
+    test_f16_to_i32_vec(vec2(1., 2.));
+    test_f16_to_u32_vec(vec2(1., 2.));
+    test_f16_to_i64_vec(vec2(1., 2.));
+    test_f16_to_u64_vec(vec2(1., 2.));
+    test_f32_to_i32_vec(vec2(1., 2.));
+    test_f32_to_u32_vec(vec2(1., 2.));
+    test_f32_to_i64_vec(vec2(1., 2.));
+    test_f32_to_u64_vec(vec2(1., 2.));
+    test_f64_to_i32_vec(vec2(1., 2.));
+    test_f64_to_u32_vec(vec2(1., 2.));
+    test_f64_to_i64_vec(vec2(1., 2.));
+    test_f64_to_u64_vec(vec2(1., 2.));
+}

naga/tests/in/wgsl/debug-symbol-large-source.wgsl

@@ -7418,6 +7418,8 @@ fn color23(p: vec2<f32>) -> vec3<f32> {
 
 @fragment
 fn fs_main(in: VertexOutput) -> @location(0) vec4<f32> {
+    color23(vec2(1, 2));
+
     var color = smoothstep(vec3<f32>(0.0), vec3<f32>(0.1), fract(in.world_pos));
     color = mix(vec3<f32>(0.5, 0.1, 0.7), vec3<f32>(0.2, 0.2, 0.2), vec3<f32>(color.x * color.y * color.z));
 

naga/tests/in/wgsl/debug-symbol-terrain.wgsl

@@ -275,6 +275,8 @@ fn color23(p: vec2<f32>) -> vec3<f32> {
 
 @fragment
 fn fs_main(in: VertexOutput) -> @location(0) vec4<f32> {
+    color23(vec2(1, 2));
+
     var color = smoothstep(vec3<f32>(0.0), vec3<f32>(0.1), fract(in.world_pos));
     color = mix(vec3<f32>(0.5, 0.1, 0.7), vec3<f32>(0.2, 0.2, 0.2), vec3<f32>(color.x * color.y * color.z));
 

naga/tests/in/wgsl/diagnostic-filter.wgsl

@@ -4,3 +4,9 @@ fn thing() {}
 
 @diagnostic(warning, derivative_uniformity)
 fn with_diagnostic() {}
+
+@compute @workgroup_size(1)
+fn main() {
+    thing();
+    with_diagnostic();
+}

naga/tests/in/wgsl/index-by-value.wgsl

@@ -12,9 +12,16 @@ fn index_let_matrix(i: i32, j: i32) -> f32 {
   return a[i][j];
 }
 
-@vertex
+fn index_let_array_1d(vi: u32) -> vec4<f32> {
-fn index_let_array_1d(@builtin(vertex_index) vi: u32) -> @builtin(position) vec4<f32> {
 	let arr = array<i32, 5>(1, 2, 3, 4, 5);
 	let value = arr[vi];
 	return vec4<f32>(vec4<i32>(value));
 }
+
+@vertex
+fn main(@builtin(vertex_index) vi: u32) -> @builtin(position) vec4<f32> {
+    index_arg_array(array(1, 2, 3, 4, 5), 6);
+    index_let_array(1, 2);
+    index_let_matrix(1, 2);
+    return index_let_array_1d(vi);
+}

naga/tests/in/wgsl/lexical-scopes.wgsl

@@ -52,3 +52,13 @@ fn switchLexicalScope(a: i32) {
     }
     let test = a == 2;
 }
+
+@compute @workgroup_size(1)
+fn main() {
+    blockLexicalScope(false);
+    ifLexicalScope(true);
+    loopLexicalScope(false);
+    forLexicalScope(1.);
+    whileLexicalScope(1);
+    switchLexicalScope(1);
+}

naga/tests/in/wgsl/local-const.wgsl

@@ -24,3 +24,8 @@ fn const_in_fn() {
     const eg = ge;
     const fg = gf;
 }
+
+@compute @workgroup_size(1)
+fn main() {
+    const_in_fn();
+}

naga/tests/in/wgsl/module-scope.wgsl

@@ -24,3 +24,10 @@ var Texture: texture_2d<f32>;
 var Sampler: sampler;
 
 alias Vec2 = vec2<f32>;
+
+@fragment
+fn main() {
+    call();
+    statement();
+    returns();
+}

naga/tests/in/wgsl/must-use.wgsl

@@ -1,6 +1,3 @@
-@compute @workgroup_size(1)
-fn main() {}
-
 @must_use
 fn use_me() -> i32 { return 10; }
 
@@ -21,3 +18,11 @@ fn use_assign_let() -> i32 {
 fn use_phony_assign() {
     _ = use_me();
 }
+
+@compute @workgroup_size(1)
+fn main() {
+    use_return();
+    use_assign_var();
+    use_assign_let();
+    use_phony_assign();
+}

naga/tests/in/wgsl/operators.wgsl

@@ -287,19 +287,6 @@ fn assignment() {
     vec0[one_i]--;
 }
 
-@compute @workgroup_size(1)
-fn main(@builtin(workgroup_id) id: vec3<u32>) {
-    builtins();
-    splat(f32(id.x), i32(id.y));
-    bool_cast(v_f32_one.xyz);
-
-    logical();
-    arithmetic();
-    bit();
-    comparison();
-    assignment();
-}
-
 fn negation_avoids_prefix_decrement() {
     let i = 1;
     let i0 = -i;
@@ -321,3 +308,20 @@ fn negation_avoids_prefix_decrement() {
     let f6 = - - -(- -f);
     let f7 = (- - - - -f);
 }
+
+@compute @workgroup_size(1)
+fn main(@builtin(workgroup_id) id: vec3<u32>) {
+    builtins();
+    splat(f32(id.x), i32(id.y));
+    splat_assignment();
+    bool_cast(v_f32_one.xyz);
+
+    logical();
+    arithmetic();
+    bit();
+    comparison();
+    assignment();
+
+    negation_avoids_prefix_decrement();
+}
+

naga/tests/in/wgsl/pointer-function-arg-restrict.wgsl

@@ -59,3 +59,20 @@ fn let_binding(a: ptr<function, array<i32, 4>>, i: u32) {
 // Runtime-sized arrays can only appear in storage buffers, while (in the base
 // language) pointers can only appear in function or private space, so there
 // is no interaction to test.
+
+@compute @workgroup_size(1)
+fn main() {
+    var vec: array<vec2<i32>, 4>;
+    var mat: array<mat2x2<f32>, 4>;
+    var arr1d: array<i32, 4>;
+    var arr2d: array<array<i32, 4>, 4>;
+    var arr3d: array<array<array<i32, 4>, 4>, 4>;
+    local_var(1);
+    mat_vec_ptrs(&vec, &mat, 1);
+    argument(&arr1d, 1);
+    argument_nested_x2(&arr2d, 1, 2);
+    argument_nested_x3(&arr3d, 1, 2);
+    index_from_self(&arr1d, 1);
+    local_var_from_arg(array(1, 2, 3, 4), 5);
+    let_binding(&arr1d, 1);
+}

naga/tests/in/wgsl/pointer-function-arg-rzsw.wgsl

@@ -59,3 +59,20 @@ fn let_binding(a: ptr<function, array<i32, 4>>, i: u32) {
 // Runtime-sized arrays can only appear in storage buffers, while (in the base
 // language) pointers can only appear in function or private space, so there
 // is no interaction to test.
+
+@compute @workgroup_size(1)
+fn main() {
+    var vec: array<vec2<i32>, 4>;
+    var mat: array<mat2x2<f32>, 4>;
+    var arr1d: array<i32, 4>;
+    var arr2d: array<array<i32, 4>, 4>;
+    var arr3d: array<array<array<i32, 4>, 4>, 4>;
+    local_var(1);
+    mat_vec_ptrs(&vec, &mat, 1);
+    argument(&arr1d, 1);
+    argument_nested_x2(&arr2d, 1, 2);
+    argument_nested_x3(&arr3d, 1, 2);
+    index_from_self(&arr1d, 1);
+    local_var_from_arg(array(1, 2, 3, 4), 5);
+    let_binding(&arr1d, 1);
+}

naga/tests/in/wgsl/pointer-function-arg.wgsl

@@ -1,6 +1,3 @@
-@compute @workgroup_size(1)
-fn main() {}
-
 fn takes_ptr(p: ptr<function, i32>) {}
 fn takes_array_ptr(p: ptr<function, array<i32, 4>>) {}
 fn takes_vec_ptr(p: ptr<function, vec2<i32>>) {}
@@ -62,3 +59,20 @@ fn let_binding(a: ptr<function, array<i32, 4>>, i: u32) {
 // Runtime-sized arrays can only appear in storage buffers, while (in the base
 // language) pointers can only appear in function or private space, so there
 // is no interaction to test.
+
+@compute @workgroup_size(1)
+fn main() {
+    var vec: array<vec2<i32>, 4>;
+    var mat: array<mat2x2<f32>, 4>;
+    var arr1d: array<i32, 4>;
+    var arr2d: array<array<i32, 4>, 4>;
+    var arr3d: array<array<array<i32, 4>, 4>, 4>;
+    local_var(1);
+    mat_vec_ptrs(&vec, &mat, 1);
+    argument(&arr1d, 1);
+    argument_nested_x2(&arr2d, 1, 2);
+    argument_nested_x3(&arr3d, 1, 2);
+    index_from_self(&arr1d, 1);
+    local_var_from_arg(array(1, 2, 3, 4), 5);
+    let_binding(&arr1d, 1);
+}

naga/tests/in/wgsl/pointers.wgsl

@@ -24,3 +24,10 @@ fn index_dynamic_array(i: i32, v: u32) {
    let val = (*p)[i];
    (*p)[i] = val + v;
 }
+
+@compute @workgroup_size(1)
+fn main() {
+    f();
+    index_unsized(1, 1);
+    index_dynamic_array(1, 1);
+}

naga/tests/in/wgsl/policy-mix.wgsl

@@ -31,3 +31,8 @@ fn mock_function(c: vec2<i32>, i: i32, l: i32) -> vec4<f32> {
           in_private[i] +
           in_function[i]);
 }
+
+@compute @workgroup_size(1)
+fn main() {
+    mock_function(vec2(1, 2), 3, 4);
+}

naga/tests/in/wgsl/type-alias.wgsl

@@ -5,6 +5,7 @@ alias Mat3 = mat3x3f;
 alias I32 = i32;
 alias F32 = f32;
 
+@compute @workgroup_size(1)
 fn main() {
     let a = FVec3(0.0, 0.0, 0.0);
     let c = FVec3(0.0);

naga/tests/out/analysis/wgsl-access.info.ron

@@ -2750,6 +2750,127 @@
             dual_source_blending: false,
             diagnostic_filter_leaf: None,
         ),
+        (
+            flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"),
+            available_stages: ("VERTEX | FRAGMENT | COMPUTE"),
+            uniformity: (
+                non_uniform_result: None,
+                requirements: (""),
+            ),
+            may_kill: false,
+            sampling_set: [],
+            global_uses: [
+                (""),
+                (""),
+                (""),
+                (""),
+                (""),
+            ],
+            expressions: [
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Value(Scalar((
+                        kind: Uint,
+                        width: 4,
+                    ))),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: Some(1),
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Value(Pointer(
+                        base: 0,
+                        space: Function,
+                    )),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Value(Scalar((
+                        kind: Float,
+                        width: 4,
+                    ))),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Value(Vector(
+                        size: Quad,
+                        scalar: (
+                            kind: Float,
+                            width: 4,
+                        ),
+                    )),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Value(Scalar((
+                        kind: Float,
+                        width: 4,
+                    ))),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Value(Vector(
+                        size: Quad,
+                        scalar: (
+                            kind: Float,
+                            width: 4,
+                        ),
+                    )),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Handle(27),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: Some(7),
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Value(Pointer(
+                        base: 27,
+                        space: Function,
+                    )),
+                ),
+            ],
+            sampling: [],
+            dual_source_blending: false,
+            diagnostic_filter_leaf: None,
+        ),
         (
             flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"),
             available_stages: ("VERTEX | FRAGMENT | COMPUTE"),
@@ -2964,6 +3085,70 @@
             dual_source_blending: false,
             diagnostic_filter_leaf: None,
         ),
+        (
+            flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"),
+            available_stages: ("VERTEX | FRAGMENT | COMPUTE"),
+            uniformity: (
+                non_uniform_result: Some(0),
+                requirements: (""),
+            ),
+            may_kill: false,
+            sampling_set: [],
+            global_uses: [
+                (""),
+                (""),
+                (""),
+                (""),
+                (""),
+            ],
+            expressions: [
+                (
+                    uniformity: (
+                        non_uniform_result: Some(0),
+                        requirements: (""),
+                    ),
+                    ref_count: 2,
+                    assignable_global: None,
+                    ty: Value(Pointer(
+                        base: 29,
+                        space: Function,
+                    )),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: Some(0),
+                        requirements: (""),
+                    ),
+                    ref_count: 0,
+                    assignable_global: None,
+                    ty: Handle(0),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: Some(2),
+                        requirements: (""),
+                    ),
+                    ref_count: 2,
+                    assignable_global: None,
+                    ty: Value(Pointer(
+                        base: 31,
+                        space: Function,
+                    )),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: Some(0),
+                        requirements: (""),
+                    ),
+                    ref_count: 0,
+                    assignable_global: None,
+                    ty: Handle(0),
+                ),
+            ],
+            sampling: [],
+            dual_source_blending: false,
+            diagnostic_filter_leaf: None,
+        ),
         (
             flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"),
             available_stages: ("VERTEX | FRAGMENT | COMPUTE"),
@@ -4537,127 +4722,6 @@
             dual_source_blending: false,
             diagnostic_filter_leaf: None,
         ),
-        (
-            flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"),
-            available_stages: ("VERTEX | FRAGMENT | COMPUTE"),
-            uniformity: (
-                non_uniform_result: None,
-                requirements: (""),
-            ),
-            may_kill: false,
-            sampling_set: [],
-            global_uses: [
-                (""),
-                (""),
-                (""),
-                (""),
-                (""),
-            ],
-            expressions: [
-                (
-                    uniformity: (
-                        non_uniform_result: None,
-                        requirements: (""),
-                    ),
-                    ref_count: 1,
-                    assignable_global: None,
-                    ty: Value(Scalar((
-                        kind: Uint,
-                        width: 4,
-                    ))),
-                ),
-                (
-                    uniformity: (
-                        non_uniform_result: Some(1),
-                        requirements: (""),
-                    ),
-                    ref_count: 1,
-                    assignable_global: None,
-                    ty: Value(Pointer(
-                        base: 0,
-                        space: Function,
-                    )),
-                ),
-                (
-                    uniformity: (
-                        non_uniform_result: None,
-                        requirements: (""),
-                    ),
-                    ref_count: 1,
-                    assignable_global: None,
-                    ty: Value(Scalar((
-                        kind: Float,
-                        width: 4,
-                    ))),
-                ),
-                (
-                    uniformity: (
-                        non_uniform_result: None,
-                        requirements: (""),
-                    ),
-                    ref_count: 1,
-                    assignable_global: None,
-                    ty: Value(Vector(
-                        size: Quad,
-                        scalar: (
-                            kind: Float,
-                            width: 4,
-                        ),
-                    )),
-                ),
-                (
-                    uniformity: (
-                        non_uniform_result: None,
-                        requirements: (""),
-                    ),
-                    ref_count: 1,
-                    assignable_global: None,
-                    ty: Value(Scalar((
-                        kind: Float,
-                        width: 4,
-                    ))),
-                ),
-                (
-                    uniformity: (
-                        non_uniform_result: None,
-                        requirements: (""),
-                    ),
-                    ref_count: 1,
-                    assignable_global: None,
-                    ty: Value(Vector(
-                        size: Quad,
-                        scalar: (
-                            kind: Float,
-                            width: 4,
-                        ),
-                    )),
-                ),
-                (
-                    uniformity: (
-                        non_uniform_result: None,
-                        requirements: (""),
-                    ),
-                    ref_count: 1,
-                    assignable_global: None,
-                    ty: Handle(27),
-                ),
-                (
-                    uniformity: (
-                        non_uniform_result: Some(7),
-                        requirements: (""),
-                    ),
-                    ref_count: 1,
-                    assignable_global: None,
-                    ty: Value(Pointer(
-                        base: 27,
-                        space: Function,
-                    )),
-                ),
-            ],
-            sampling: [],
-            dual_source_blending: false,
-            diagnostic_filter_leaf: None,
-        ),
         (
             flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"),
             available_stages: ("VERTEX | FRAGMENT | COMPUTE"),
@@ -4677,45 +4741,51 @@
             expressions: [
                 (
                     uniformity: (
-                        non_uniform_result: Some(0),
+                        non_uniform_result: None,
                         requirements: (""),
                     ),
-                    ref_count: 2,
+                    ref_count: 1,
                     assignable_global: None,
-                    ty: Value(Pointer(
+                    ty: Value(Scalar((
-                        base: 29,
+                        kind: Bool,
-                        space: Function,
+                        width: 1,
-                    )),
+                    ))),
-                ),
-                (
-                    uniformity: (
-                        non_uniform_result: Some(0),
-                        requirements: (""),
-                    ),
-                    ref_count: 0,
-                    assignable_global: None,
-                    ty: Handle(0),
                 ),
                 (
                     uniformity: (
                         non_uniform_result: Some(2),
                         requirements: (""),
                     ),
-                    ref_count: 2,
+                    ref_count: 0,
                     assignable_global: None,
-                    ty: Value(Pointer(
+                    ty: Handle(33),
-                        base: 31,
-                        space: Function,
-                    )),
                 ),
                 (
                     uniformity: (
-                        non_uniform_result: Some(0),
+                        non_uniform_result: Some(2),
                         requirements: (""),
                     ),
                     ref_count: 0,
                     assignable_global: None,
-                    ty: Handle(0),
+                    ty: Handle(2),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 0,
+                    assignable_global: None,
+                    ty: Handle(2),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: Some(1),
+                        requirements: (""),
+                    ),
+                    ref_count: 0,
+                    assignable_global: None,
+                    ty: Handle(2),
                 ),
             ],
             sampling: [],

naga/tests/out/glsl/wgsl-abstract-types-builtins.f.Compute.glsl

@@ -0,0 +1,67 @@
+#version 310 es
+
+precision highp float;
+precision highp int;
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+
+
+void main() {
+    int clamp_aiaiai = 1;
+    float clamp_aiaiaf = 1.0;
+    int clamp_aiaii = 1;
+    float clamp_aiaif = 1.0;
+    float clamp_aiafai = 1.0;
+    float clamp_aiafaf = 1.0;
+    float clamp_aiaff = 1.0;
+    int clamp_aiiai = 1;
+    int clamp_aiii = 1;
+    float clamp_aifai = 1.0;
+    float clamp_aifaf = 1.0;
+    float clamp_aiff = 1.0;
+    float clamp_afaiai = 1.0;
+    float clamp_afaiaf = 1.0;
+    float clamp_afaif = 1.0;
+    float clamp_afafai = 1.0;
+    float clamp_afafaf = 1.0;
+    float clamp_afaff = 1.0;
+    float clamp_affai = 1.0;
+    float clamp_affaf = 1.0;
+    float clamp_afff = 1.0;
+    int clamp_iaiai = 1;
+    int clamp_iaii = 1;
+    int clamp_iiai = 1;
+    int clamp_iii = 1;
+    float clamp_faiai = 1.0;
+    float clamp_faiaf = 1.0;
+    float clamp_faif = 1.0;
+    float clamp_fafai = 1.0;
+    float clamp_fafaf = 1.0;
+    float clamp_faff = 1.0;
+    float clamp_ffai = 1.0;
+    float clamp_ffaf = 1.0;
+    float clamp_fff = 1.0;
+    int min_aiai = 1;
+    float min_aiaf = 1.0;
+    int min_aii = 1;
+    float min_aif = 1.0;
+    float min_afai = 1.0;
+    float min_afaf = 1.0;
+    float min_aff = 1.0;
+    int min_iai = 1;
+    int min_ii = 1;
+    float min_fai = 1.0;
+    float min_faf = 1.0;
+    float min_ff = 1.0;
+    float pow_aiai = 1.0;
+    float pow_aiaf = 1.0;
+    float pow_aif = 1.0;
+    float pow_afai = 1.0;
+    float pow_afaf = 1.0;
+    float pow_aff = 1.0;
+    float pow_fai = 1.0;
+    float pow_faf = 1.0;
+    float pow_ff = 1.0;
+    return;
+}
+

naga/tests/out/glsl/wgsl-abstract-types-function-calls.main.Compute.glsl

@@ -0,0 +1,87 @@
+#version 310 es
+
+precision highp float;
+precision highp int;
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+
+
+void func_f(float a) {
+    return;
+}
+
+void func_i(int a_1) {
+    return;
+}
+
+void func_u(uint a_2) {
+    return;
+}
+
+void func_vf(vec2 a_3) {
+    return;
+}
+
+void func_vi(ivec2 a_4) {
+    return;
+}
+
+void func_vu(uvec2 a_5) {
+    return;
+}
+
+void func_mf(mat2x2 a_6) {
+    return;
+}
+
+void func_af(float a_7[2]) {
+    return;
+}
+
+void func_ai(int a_8[2]) {
+    return;
+}
+
+void func_au(uint a_9[2]) {
+    return;
+}
+
+void func_f_i(float a_10, int b) {
+    return;
+}
+
+void main() {
+    func_f(0.0);
+    func_f(0.0);
+    func_i(0);
+    func_u(0u);
+    func_f(0.0);
+    func_f(0.0);
+    func_i(0);
+    func_u(0u);
+    func_vf(vec2(0.0));
+    func_vf(vec2(0.0));
+    func_vi(ivec2(0));
+    func_vu(uvec2(0u));
+    func_vf(vec2(0.0));
+    func_vf(vec2(0.0));
+    func_vi(ivec2(0));
+    func_vu(uvec2(0u));
+    func_mf(mat2x2(vec2(0.0), vec2(0.0)));
+    func_mf(mat2x2(vec2(0.0), vec2(0.0)));
+    func_mf(mat2x2(vec2(0.0), vec2(0.0)));
+    func_af(float[2](0.0, 0.0));
+    func_af(float[2](0.0, 0.0));
+    func_ai(int[2](0, 0));
+    func_au(uint[2](0u, 0u));
+    func_af(float[2](0.0, 0.0));
+    func_af(float[2](0.0, 0.0));
+    func_ai(int[2](0, 0));
+    func_au(uint[2](0u, 0u));
+    func_f_i(0.0, 0);
+    func_f_i(0.0, 0);
+    func_f_i(0.0, 0);
+    func_f_i(0.0, 0);
+    return;
+}
+

naga/tests/out/glsl/wgsl-abstract-types-let.main.Compute.glsl

@@ -0,0 +1,125 @@
+#version 310 es
+
+precision highp float;
+precision highp int;
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+
+
+void all_constant_arguments() {
+    ivec2 xvipaiai = ivec2(42, 43);
+    uvec2 xvupaiai = uvec2(44u, 45u);
+    vec2 xvfpaiai = vec2(46.0, 47.0);
+    vec2 xvfpafaf = vec2(48.0, 49.0);
+    vec2 xvfpaiaf = vec2(48.0, 49.0);
+    uvec2 xvupuai = uvec2(42u, 43u);
+    uvec2 xvupaiu = uvec2(42u, 43u);
+    uvec2 xvuuai = uvec2(42u, 43u);
+    uvec2 xvuaiu = uvec2(42u, 43u);
+    ivec2 xvip = ivec2(0, 0);
+    uvec2 xvup = uvec2(0u, 0u);
+    vec2 xvfp = vec2(0.0, 0.0);
+    mat2x2 xmfp = mat2x2(vec2(0.0, 0.0), vec2(0.0, 0.0));
+    mat2x2 xmfpaiaiaiai = mat2x2(vec2(1.0, 2.0), vec2(3.0, 4.0));
+    mat2x2 xmfpafaiaiai = mat2x2(vec2(1.0, 2.0), vec2(3.0, 4.0));
+    mat2x2 xmfpaiafaiai = mat2x2(vec2(1.0, 2.0), vec2(3.0, 4.0));
+    mat2x2 xmfpaiaiafai = mat2x2(vec2(1.0, 2.0), vec2(3.0, 4.0));
+    mat2x2 xmfpaiaiaiaf = mat2x2(vec2(1.0, 2.0), vec2(3.0, 4.0));
+    mat2x2 xmfp_faiaiai = mat2x2(vec2(1.0, 2.0), vec2(3.0, 4.0));
+    mat2x2 xmfpai_faiai = mat2x2(vec2(1.0, 2.0), vec2(3.0, 4.0));
+    mat2x2 xmfpaiai_fai = mat2x2(vec2(1.0, 2.0), vec2(3.0, 4.0));
+    mat2x2 xmfpaiaiai_f = mat2x2(vec2(1.0, 2.0), vec2(3.0, 4.0));
+    ivec2 xvispai = ivec2(1);
+    vec2 xvfspaf = vec2(1.0);
+    ivec2 xvis_ai = ivec2(1);
+    uvec2 xvus_ai = uvec2(1u);
+    vec2 xvfs_ai = vec2(1.0);
+    vec2 xvfs_af = vec2(1.0);
+    float xafafaf[2] = float[2](1.0, 2.0);
+    float xaf_faf[2] = float[2](1.0, 2.0);
+    float xafaf_f[2] = float[2](1.0, 2.0);
+    float xafaiai[2] = float[2](1.0, 2.0);
+    int xai_iai[2] = int[2](1, 2);
+    int xaiai_i[2] = int[2](1, 2);
+    int xaipaiai[2] = int[2](1, 2);
+    float xafpaiai[2] = float[2](1.0, 2.0);
+    float xafpaiaf[2] = float[2](1.0, 2.0);
+    float xafpafai[2] = float[2](1.0, 2.0);
+    float xafpafaf[2] = float[2](1.0, 2.0);
+    ivec3 xavipai[1] = ivec3[1](ivec3(1));
+    vec3 xavfpai[1] = vec3[1](vec3(1.0));
+    vec3 xavfpaf[1] = vec3[1](vec3(1.0));
+    ivec2 xvisai = ivec2(1);
+    uvec2 xvusai = uvec2(1u);
+    vec2 xvfsai = vec2(1.0);
+    vec2 xvfsaf = vec2(1.0);
+    int iaipaiai[2] = int[2](1, 2);
+    float iafpaiaf[2] = float[2](1.0, 2.0);
+    float iafpafai[2] = float[2](1.0, 2.0);
+    float iafpafaf[2] = float[2](1.0, 2.0);
+    return;
+}
+
+void mixed_constant_and_runtime_arguments() {
+    uint u = 0u;
+    int i = 0;
+    float f = 0.0;
+    uint _e3 = u;
+    uvec2 xvupuai_1 = uvec2(_e3, 43u);
+    uint _e6 = u;
+    uvec2 xvupaiu_1 = uvec2(42u, _e6);
+    float _e9 = f;
+    vec2 xvfpfai = vec2(_e9, 47.0);
+    float _e12 = f;
+    vec2 xvfpfaf = vec2(_e12, 49.0);
+    uint _e15 = u;
+    uvec2 xvuuai_1 = uvec2(_e15, 43u);
+    uint _e18 = u;
+    uvec2 xvuaiu_1 = uvec2(42u, _e18);
+    float _e21 = f;
+    mat2x2 xmfp_faiaiai_1 = mat2x2(vec2(_e21, 2.0), vec2(3.0, 4.0));
+    float _e28 = f;
+    mat2x2 xmfpai_faiai_1 = mat2x2(vec2(1.0, _e28), vec2(3.0, 4.0));
+    float _e35 = f;
+    mat2x2 xmfpaiai_fai_1 = mat2x2(vec2(1.0, 2.0), vec2(_e35, 4.0));
+    float _e42 = f;
+    mat2x2 xmfpaiaiai_f_1 = mat2x2(vec2(1.0, 2.0), vec2(3.0, _e42));
+    float _e49 = f;
+    float xaf_faf_1[2] = float[2](_e49, 2.0);
+    float _e52 = f;
+    float xafaf_f_1[2] = float[2](1.0, _e52);
+    float _e55 = f;
+    float xaf_fai[2] = float[2](_e55, 2.0);
+    float _e58 = f;
+    float xafai_f[2] = float[2](1.0, _e58);
+    int _e61 = i;
+    int xai_iai_1[2] = int[2](_e61, 2);
+    int _e64 = i;
+    int xaiai_i_1[2] = int[2](1, _e64);
+    float _e67 = f;
+    float xafp_faf[2] = float[2](_e67, 2.0);
+    float _e70 = f;
+    float xafpaf_f[2] = float[2](1.0, _e70);
+    float _e73 = f;
+    float xafp_fai[2] = float[2](_e73, 2.0);
+    float _e76 = f;
+    float xafpai_f[2] = float[2](1.0, _e76);
+    int _e79 = i;
+    int xaip_iai[2] = int[2](_e79, 2);
+    int _e82 = i;
+    int xaipai_i[2] = int[2](1, _e82);
+    int _e85 = i;
+    ivec2 xvisi = ivec2(_e85);
+    uint _e87 = u;
+    uvec2 xvusu = uvec2(_e87);
+    float _e89 = f;
+    vec2 xvfsf = vec2(_e89);
+    return;
+}
+
+void main() {
+    all_constant_arguments();
+    mixed_constant_and_runtime_arguments();
+    return;
+}
+

naga/tests/out/glsl/wgsl-abstract-types-operators.main.Compute.glsl

@@ -0,0 +1,117 @@
+#version 310 es
+
+precision highp float;
+precision highp int;
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+
+const float plus_fafaf_1 = 3.0;
+const float plus_fafai_1 = 3.0;
+const float plus_faf_f_1 = 3.0;
+const float plus_faiaf_1 = 3.0;
+const float plus_faiai_1 = 3.0;
+const float plus_fai_f_1 = 3.0;
+const float plus_f_faf_1 = 3.0;
+const float plus_f_fai_1 = 3.0;
+const float plus_f_f_f_1 = 3.0;
+const int plus_iaiai_1 = 3;
+const int plus_iai_i_1 = 3;
+const int plus_i_iai_1 = 3;
+const int plus_i_i_i_1 = 3;
+const uint plus_uaiai_1 = 3u;
+const uint plus_uai_u_1 = 3u;
+const uint plus_u_uai_1 = 3u;
+const uint plus_u_u_u_1 = 3u;
+const uint bitflip_u_u = 0u;
+const uint bitflip_uai = 0u;
+const int least_i32_ = -2147483648;
+const float least_f32_ = -3.4028235e38;
+const int shl_iaiai = 4;
+const int shl_iai_u_1 = 4;
+const uint shl_uaiai = 4u;
+const uint shl_uai_u = 4u;
+const int shr_iaiai = 0;
+const int shr_iai_u_1 = 0;
+const uint shr_uaiai = 0u;
+const uint shr_uai_u = 0u;
+const int wgpu_4492_ = -2147483648;
+
+shared uint a[64];
+
+
+void runtime_values() {
+    float f = 42.0;
+    int i = 43;
+    uint u = 44u;
+    float plus_fafaf = 3.0;
+    float plus_fafai = 3.0;
+    float plus_faf_f = 0.0;
+    float plus_faiaf = 3.0;
+    float plus_faiai = 3.0;
+    float plus_fai_f = 0.0;
+    float plus_f_faf = 0.0;
+    float plus_f_fai = 0.0;
+    float plus_f_f_f = 0.0;
+    int plus_iaiai = 3;
+    int plus_iai_i = 0;
+    int plus_i_iai = 0;
+    int plus_i_i_i = 0;
+    uint plus_uaiai = 3u;
+    uint plus_uai_u = 0u;
+    uint plus_u_uai = 0u;
+    uint plus_u_u_u = 0u;
+    int shl_iai_u = 0;
+    int shr_iai_u = 0;
+    float _e8 = f;
+    plus_faf_f = (1.0 + _e8);
+    float _e14 = f;
+    plus_fai_f = (1.0 + _e14);
+    float _e18 = f;
+    plus_f_faf = (_e18 + 2.0);
+    float _e22 = f;
+    plus_f_fai = (_e22 + 2.0);
+    float _e26 = f;
+    float _e27 = f;
+    plus_f_f_f = (_e26 + _e27);
+    int _e31 = i;
+    plus_iai_i = (1 + _e31);
+    int _e35 = i;
+    plus_i_iai = (_e35 + 2);
+    int _e39 = i;
+    int _e40 = i;
+    plus_i_i_i = (_e39 + _e40);
+    uint _e44 = u;
+    plus_uai_u = (1u + _e44);
+    uint _e48 = u;
+    plus_u_uai = (_e48 + 2u);
+    uint _e52 = u;
+    uint _e53 = u;
+    plus_u_u_u = (_e52 + _e53);
+    uint _e56 = u;
+    shl_iai_u = (1 << _e56);
+    uint _e60 = u;
+    shr_iai_u = (1 << _e60);
+    return;
+}
+
+void wgpu_4445_() {
+    return;
+}
+
+void wgpu_4435_() {
+    uint y = a[(1 - 1)];
+    return;
+}
+
+void main() {
+    if (gl_LocalInvocationID == uvec3(0u)) {
+        a = uint[64](0u, 0u, 0u, 0u, 0u, 0u, 0u, 0u, 0u, 0u, 0u, 0u, 0u, 0u, 0u, 0u, 0u, 0u, 0u, 0u, 0u, 0u, 0u, 0u, 0u, 0u, 0u, 0u, 0u, 0u, 0u, 0u, 0u, 0u, 0u, 0u, 0u, 0u, 0u, 0u, 0u, 0u, 0u, 0u, 0u, 0u, 0u, 0u, 0u, 0u, 0u, 0u, 0u, 0u, 0u, 0u, 0u, 0u, 0u, 0u, 0u, 0u, 0u, 0u);
+    }
+    memoryBarrierShared();
+    barrier();
+    runtime_values();
+    wgpu_4445_();
+    wgpu_4435_();
+    return;
+}
+

naga/tests/out/glsl/wgsl-abstract-types-return.main.Compute.glsl

@@ -39,6 +39,14 @@ vec2 return_vec2f32_const_ai() {
 }
 
 void main() {
+    int _e0 = return_i32_ai();
+    uint _e1 = return_u32_ai();
+    float _e2 = return_f32_ai();
+    float _e3 = return_f32_af();
+    vec2 _e4 = return_vec2f32_ai();
+    float _e5[4] = return_arrf32_ai();
+    float _e6 = return_const_f32_const_ai();
+    vec2 _e7 = return_vec2f32_const_ai();
     return;
 }
 

naga/tests/out/glsl/wgsl-abstract-types-var.main.Compute.glsl

@@ -0,0 +1,246 @@
+#version 310 es
+
+precision highp float;
+precision highp int;
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+
+
+void all_constant_arguments() {
+    ivec2 xvipaiai = ivec2(42, 43);
+    uvec2 xvupaiai = uvec2(44u, 45u);
+    vec2 xvfpaiai = vec2(46.0, 47.0);
+    vec2 xvfpafaf = vec2(48.0, 49.0);
+    vec2 xvfpaiaf = vec2(48.0, 49.0);
+    uvec2 xvupuai = uvec2(42u, 43u);
+    uvec2 xvupaiu = uvec2(42u, 43u);
+    uvec2 xvuuai = uvec2(42u, 43u);
+    uvec2 xvuaiu = uvec2(42u, 43u);
+    ivec2 xvip = ivec2(0, 0);
+    uvec2 xvup = uvec2(0u, 0u);
+    vec2 xvfp = vec2(0.0, 0.0);
+    mat2x2 xmfp = mat2x2(vec2(0.0, 0.0), vec2(0.0, 0.0));
+    mat2x2 xmfpaiaiaiai = mat2x2(vec2(1.0, 2.0), vec2(3.0, 4.0));
+    mat2x2 xmfpafaiaiai = mat2x2(vec2(1.0, 2.0), vec2(3.0, 4.0));
+    mat2x2 xmfpaiafaiai = mat2x2(vec2(1.0, 2.0), vec2(3.0, 4.0));
+    mat2x2 xmfpaiaiafai = mat2x2(vec2(1.0, 2.0), vec2(3.0, 4.0));
+    mat2x2 xmfpaiaiaiaf = mat2x2(vec2(1.0, 2.0), vec2(3.0, 4.0));
+    mat2x2 xmfp_faiaiai = mat2x2(vec2(1.0, 2.0), vec2(3.0, 4.0));
+    mat2x2 xmfpai_faiai = mat2x2(vec2(1.0, 2.0), vec2(3.0, 4.0));
+    mat2x2 xmfpaiai_fai = mat2x2(vec2(1.0, 2.0), vec2(3.0, 4.0));
+    mat2x2 xmfpaiaiai_f = mat2x2(vec2(1.0, 2.0), vec2(3.0, 4.0));
+    ivec2 xvispai = ivec2(1);
+    vec2 xvfspaf = vec2(1.0);
+    ivec2 xvis_ai = ivec2(1);
+    uvec2 xvus_ai = uvec2(1u);
+    vec2 xvfs_ai = vec2(1.0);
+    vec2 xvfs_af = vec2(1.0);
+    float xafafaf[2] = float[2](1.0, 2.0);
+    float xaf_faf[2] = float[2](1.0, 2.0);
+    float xafaf_f[2] = float[2](1.0, 2.0);
+    float xafaiai[2] = float[2](1.0, 2.0);
+    int xai_iai[2] = int[2](1, 2);
+    int xaiai_i[2] = int[2](1, 2);
+    int xaipaiai[2] = int[2](1, 2);
+    float xafpaiai[2] = float[2](1.0, 2.0);
+    float xafpaiaf[2] = float[2](1.0, 2.0);
+    float xafpafai[2] = float[2](1.0, 2.0);
+    float xafpafaf[2] = float[2](1.0, 2.0);
+    ivec3 xavipai[1] = ivec3[1](ivec3(1));
+    vec3 xavfpai[1] = vec3[1](vec3(1.0));
+    vec3 xavfpaf[1] = vec3[1](vec3(1.0));
+    ivec2 xvisai = ivec2(1);
+    uvec2 xvusai = uvec2(1u);
+    vec2 xvfsai = vec2(1.0);
+    vec2 xvfsaf = vec2(1.0);
+    int iaipaiai[2] = int[2](1, 2);
+    float iafpaiaf[2] = float[2](1.0, 2.0);
+    float iafpafai[2] = float[2](1.0, 2.0);
+    float iafpafaf[2] = float[2](1.0, 2.0);
+    xvipaiai = ivec2(42, 43);
+    xvupaiai = uvec2(44u, 45u);
+    xvfpaiai = vec2(46.0, 47.0);
+    xvfpafaf = vec2(48.0, 49.0);
+    xvfpaiaf = vec2(48.0, 49.0);
+    xvupuai = uvec2(42u, 43u);
+    xvupaiu = uvec2(42u, 43u);
+    xvuuai = uvec2(42u, 43u);
+    xvuaiu = uvec2(42u, 43u);
+    xvip = ivec2(0, 0);
+    xvup = uvec2(0u, 0u);
+    xvfp = vec2(0.0, 0.0);
+    xmfp = mat2x2(vec2(0.0, 0.0), vec2(0.0, 0.0));
+    xmfpaiaiaiai = mat2x2(vec2(1.0, 2.0), vec2(3.0, 4.0));
+    xmfpafaiaiai = mat2x2(vec2(1.0, 2.0), vec2(3.0, 4.0));
+    xmfpaiafaiai = mat2x2(vec2(1.0, 2.0), vec2(3.0, 4.0));
+    xmfpaiaiafai = mat2x2(vec2(1.0, 2.0), vec2(3.0, 4.0));
+    xmfpaiaiaiaf = mat2x2(vec2(1.0, 2.0), vec2(3.0, 4.0));
+    xmfp_faiaiai = mat2x2(vec2(1.0, 2.0), vec2(3.0, 4.0));
+    xmfpai_faiai = mat2x2(vec2(1.0, 2.0), vec2(3.0, 4.0));
+    xmfpaiai_fai = mat2x2(vec2(1.0, 2.0), vec2(3.0, 4.0));
+    xmfpaiaiai_f = mat2x2(vec2(1.0, 2.0), vec2(3.0, 4.0));
+    xvispai = ivec2(1);
+    xvfspaf = vec2(1.0);
+    xvis_ai = ivec2(1);
+    xvus_ai = uvec2(1u);
+    xvfs_ai = vec2(1.0);
+    xvfs_af = vec2(1.0);
+    xafafaf = float[2](1.0, 2.0);
+    xaf_faf = float[2](1.0, 2.0);
+    xafaf_f = float[2](1.0, 2.0);
+    xafaiai = float[2](1.0, 2.0);
+    xai_iai = int[2](1, 2);
+    xaiai_i = int[2](1, 2);
+    xaipaiai = int[2](1, 2);
+    xafpaiai = float[2](1.0, 2.0);
+    xafpaiaf = float[2](1.0, 2.0);
+    xafpafai = float[2](1.0, 2.0);
+    xafpafaf = float[2](1.0, 2.0);
+    xavipai = ivec3[1](ivec3(1));
+    xavfpai = vec3[1](vec3(1.0));
+    xavfpaf = vec3[1](vec3(1.0));
+    xvisai = ivec2(1);
+    xvusai = uvec2(1u);
+    xvfsai = vec2(1.0);
+    xvfsaf = vec2(1.0);
+    iaipaiai = int[2](1, 2);
+    iafpaiaf = float[2](1.0, 2.0);
+    iafpafai = float[2](1.0, 2.0);
+    iafpafaf = float[2](1.0, 2.0);
+    return;
+}
+
+void mixed_constant_and_runtime_arguments() {
+    uint u = 0u;
+    int i = 0;
+    float f = 0.0;
+    uvec2 xvupuai_1 = uvec2(0u);
+    uvec2 xvupaiu_1 = uvec2(0u);
+    vec2 xvfpfai = vec2(0.0);
+    vec2 xvfpfaf = vec2(0.0);
+    uvec2 xvuuai_1 = uvec2(0u);
+    uvec2 xvuaiu_1 = uvec2(0u);
+    mat2x2 xmfp_faiaiai_1 = mat2x2(0.0);
+    mat2x2 xmfpai_faiai_1 = mat2x2(0.0);
+    mat2x2 xmfpaiai_fai_1 = mat2x2(0.0);
+    mat2x2 xmfpaiaiai_f_1 = mat2x2(0.0);
+    float xaf_faf_1[2] = float[2](0.0, 0.0);
+    float xafaf_f_1[2] = float[2](0.0, 0.0);
+    float xaf_fai[2] = float[2](0.0, 0.0);
+    float xafai_f[2] = float[2](0.0, 0.0);
+    int xai_iai_1[2] = int[2](0, 0);
+    int xaiai_i_1[2] = int[2](0, 0);
+    float xafp_faf[2] = float[2](0.0, 0.0);
+    float xafpaf_f[2] = float[2](0.0, 0.0);
+    float xafp_fai[2] = float[2](0.0, 0.0);
+    float xafpai_f[2] = float[2](0.0, 0.0);
+    int xaip_iai[2] = int[2](0, 0);
+    int xaipai_i[2] = int[2](0, 0);
+    ivec2 xvisi = ivec2(0);
+    uvec2 xvusu = uvec2(0u);
+    vec2 xvfsf = vec2(0.0);
+    uint _e3 = u;
+    xvupuai_1 = uvec2(_e3, 43u);
+    uint _e7 = u;
+    xvupaiu_1 = uvec2(42u, _e7);
+    float _e11 = f;
+    xvfpfai = vec2(_e11, 47.0);
+    float _e15 = f;
+    xvfpfaf = vec2(_e15, 49.0);
+    uint _e19 = u;
+    xvuuai_1 = uvec2(_e19, 43u);
+    uint _e23 = u;
+    xvuaiu_1 = uvec2(42u, _e23);
+    float _e27 = f;
+    xmfp_faiaiai_1 = mat2x2(vec2(_e27, 2.0), vec2(3.0, 4.0));
+    float _e35 = f;
+    xmfpai_faiai_1 = mat2x2(vec2(1.0, _e35), vec2(3.0, 4.0));
+    float _e43 = f;
+    xmfpaiai_fai_1 = mat2x2(vec2(1.0, 2.0), vec2(_e43, 4.0));
+    float _e51 = f;
+    xmfpaiaiai_f_1 = mat2x2(vec2(1.0, 2.0), vec2(3.0, _e51));
+    float _e59 = f;
+    xaf_faf_1 = float[2](_e59, 2.0);
+    float _e63 = f;
+    xafaf_f_1 = float[2](1.0, _e63);
+    float _e67 = f;
+    xaf_fai = float[2](_e67, 2.0);
+    float _e71 = f;
+    xafai_f = float[2](1.0, _e71);
+    int _e75 = i;
+    xai_iai_1 = int[2](_e75, 2);
+    int _e79 = i;
+    xaiai_i_1 = int[2](1, _e79);
+    float _e83 = f;
+    xafp_faf = float[2](_e83, 2.0);
+    float _e87 = f;
+    xafpaf_f = float[2](1.0, _e87);
+    float _e91 = f;
+    xafp_fai = float[2](_e91, 2.0);
+    float _e95 = f;
+    xafpai_f = float[2](1.0, _e95);
+    int _e99 = i;
+    xaip_iai = int[2](_e99, 2);
+    int _e103 = i;
+    xaipai_i = int[2](1, _e103);
+    int _e107 = i;
+    xvisi = ivec2(_e107);
+    uint _e110 = u;
+    xvusu = uvec2(_e110);
+    float _e113 = f;
+    xvfsf = vec2(_e113);
+    uint _e116 = u;
+    xvupuai_1 = uvec2(_e116, 43u);
+    uint _e119 = u;
+    xvupaiu_1 = uvec2(42u, _e119);
+    uint _e122 = u;
+    xvuuai_1 = uvec2(_e122, 43u);
+    uint _e125 = u;
+    xvuaiu_1 = uvec2(42u, _e125);
+    float _e128 = f;
+    xmfp_faiaiai_1 = mat2x2(vec2(_e128, 2.0), vec2(3.0, 4.0));
+    float _e135 = f;
+    xmfpai_faiai_1 = mat2x2(vec2(1.0, _e135), vec2(3.0, 4.0));
+    float _e142 = f;
+    xmfpaiai_fai_1 = mat2x2(vec2(1.0, 2.0), vec2(_e142, 4.0));
+    float _e149 = f;
+    xmfpaiaiai_f_1 = mat2x2(vec2(1.0, 2.0), vec2(3.0, _e149));
+    float _e156 = f;
+    xaf_faf_1 = float[2](_e156, 2.0);
+    float _e159 = f;
+    xafaf_f_1 = float[2](1.0, _e159);
+    float _e162 = f;
+    xaf_fai = float[2](_e162, 2.0);
+    float _e165 = f;
+    xafai_f = float[2](1.0, _e165);
+    int _e168 = i;
+    xai_iai_1 = int[2](_e168, 2);
+    int _e171 = i;
+    xaiai_i_1 = int[2](1, _e171);
+    float _e174 = f;
+    xafp_faf = float[2](_e174, 2.0);
+    float _e177 = f;
+    xafpaf_f = float[2](1.0, _e177);
+    float _e180 = f;
+    xafp_fai = float[2](_e180, 2.0);
+    float _e183 = f;
+    xafpai_f = float[2](1.0, _e183);
+    int _e186 = i;
+    xaip_iai = int[2](_e186, 2);
+    int _e189 = i;
+    xaipai_i = int[2](1, _e189);
+    int _e192 = i;
+    xvisi = ivec2(_e192);
+    uint _e194 = u;
+    xvusu = uvec2(_e194);
+    float _e196 = f;
+    xvfsf = vec2(_e196);
+    return;
+}
+
+void main() {
+    all_constant_arguments();
+    mixed_constant_and_runtime_arguments();
+    return;
+}
+

naga/tests/out/glsl/wgsl-access.assign_through_ptr.Compute.glsl

@@ -1,119 +0,0 @@
-#version 310 es
-
-precision highp float;
-precision highp int;
-
-layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
-
-struct GlobalConst {
-    uint a;
-    uvec3 b;
-    int c;
-};
-struct AlignedWrapper {
-    int value;
-};
-struct Baz {
-    mat3x2 m;
-};
-struct MatCx2InArray {
-    mat4x2 am[2];
-};
-struct AssignToMember {
-    uint x;
-};
-struct S {
-    int m;
-};
-struct Inner {
-    int delicious;
-};
-struct Outer {
-    Inner om_nom_nom;
-    uint thing;
-};
-
-float read_from_private(inout float foo_1) {
-    float _e1 = foo_1;
-    return _e1;
-}
-
-float test_arr_as_arg(float a[5][10]) {
-    return a[4][9];
-}
-
-void assign_through_ptr_fn(inout uint p) {
-    p = 42u;
-    return;
-}
-
-void assign_array_through_ptr_fn(inout vec4 foo_2[2]) {
-    foo_2 = vec4[2](vec4(1.0), vec4(2.0));
-    return;
-}
-
-uint fetch_arg_ptr_member(inout AssignToMember p_1) {
-    uint _e2 = p_1.x;
-    return _e2;
-}
-
-void assign_to_arg_ptr_member(inout AssignToMember p_2) {
-    p_2.x = 10u;
-    return;
-}
-
-uint fetch_arg_ptr_array_element(inout uint p_3[4]) {
-    uint _e2 = p_3[1];
-    return _e2;
-}
-
-void assign_to_arg_ptr_array_element(inout uint p_4[4]) {
-    p_4[1] = 10u;
-    return;
-}
-
-bool index_ptr(bool value) {
-    bool a_1[1] = bool[1](false);
-    a_1 = bool[1](value);
-    bool _e4 = a_1[0];
-    return _e4;
-}
-
-int member_ptr() {
-    S s = S(42);
-    int _e4 = s.m;
-    return _e4;
-}
-
-int let_members_of_members() {
-    Inner inner_1 = Outer(Inner(0), 0u).om_nom_nom;
-    int delishus_1 = inner_1.delicious;
-    if ((Outer(Inner(0), 0u).thing != uint(delishus_1))) {
-    }
-    return Outer(Inner(0), 0u).om_nom_nom.delicious;
-}
-
-int var_members_of_members() {
-    Outer thing = Outer(Inner(0), 0u);
-    Inner inner = Inner(0);
-    int delishus = 0;
-    Inner _e3 = thing.om_nom_nom;
-    inner = _e3;
-    int _e6 = inner.delicious;
-    delishus = _e6;
-    uint _e9 = thing.thing;
-    int _e10 = delishus;
-    if ((_e9 != uint(_e10))) {
-    }
-    int _e15 = thing.om_nom_nom.delicious;
-    return _e15;
-}
-
-void main() {
-    uint val = 33u;
-    vec4 arr[2] = vec4[2](vec4(6.0), vec4(7.0));
-    assign_through_ptr_fn(val);
-    assign_array_through_ptr_fn(arr);
-    return;
-}
-

naga/tests/out/glsl/wgsl-access.foo_compute.Compute.glsl

@@ -52,6 +52,14 @@ void assign_array_through_ptr_fn(inout vec4 foo_2[2]) {
     return;
 }
 
+void assign_through_ptr() {
+    uint val = 33u;
+    vec4 arr[2] = vec4[2](vec4(6.0), vec4(7.0));
+    assign_through_ptr_fn(val);
+    assign_array_through_ptr_fn(arr);
+    return;
+}
+
 uint fetch_arg_ptr_member(inout AssignToMember p_1) {
     uint _e2 = p_1.x;
     return _e2;
@@ -72,6 +80,16 @@ void assign_to_arg_ptr_array_element(inout uint p_4[4]) {
     return;
 }
 
+void assign_to_ptr_components() {
+    AssignToMember s1_ = AssignToMember(0u);
+    uint a1_[4] = uint[4](0u, 0u, 0u, 0u);
+    assign_to_arg_ptr_member(s1_);
+    uint _e1 = fetch_arg_ptr_member(s1_);
+    assign_to_arg_ptr_array_element(a1_);
+    uint _e3 = fetch_arg_ptr_array_element(a1_);
+    return;
+}
+
 bool index_ptr(bool value) {
     bool a_1[1] = bool[1](false);
     a_1 = bool[1](value);
@@ -110,12 +128,12 @@ int var_members_of_members() {
 }
 
 void main() {
-    AssignToMember s1_ = AssignToMember(0u);
+    assign_through_ptr();
-    uint a1_[4] = uint[4](0u, 0u, 0u, 0u);
+    assign_to_ptr_components();
-    assign_to_arg_ptr_member(s1_);
+    bool _e1 = index_ptr(true);
-    uint _e1 = fetch_arg_ptr_member(s1_);
+    int _e2 = member_ptr();
-    assign_to_arg_ptr_array_element(a1_);
+    int _e3 = let_members_of_members();
-    uint _e3 = fetch_arg_ptr_array_element(a1_);
+    int _e4 = var_members_of_members();
     return;
 }
 

naga/tests/out/glsl/wgsl-access.foo_frag.Fragment.glsl

@@ -62,6 +62,14 @@ void assign_array_through_ptr_fn(inout vec4 foo_2[2]) {
     return;
 }
 
+void assign_through_ptr() {
+    uint val = 33u;
+    vec4 arr[2] = vec4[2](vec4(6.0), vec4(7.0));
+    assign_through_ptr_fn(val);
+    assign_array_through_ptr_fn(arr);
+    return;
+}
+
 uint fetch_arg_ptr_member(inout AssignToMember p_1) {
     uint _e2 = p_1.x;
     return _e2;
@@ -82,6 +90,16 @@ void assign_to_arg_ptr_array_element(inout uint p_4[4]) {
     return;
 }
 
+void assign_to_ptr_components() {
+    AssignToMember s1_ = AssignToMember(0u);
+    uint a1_[4] = uint[4](0u, 0u, 0u, 0u);
+    assign_to_arg_ptr_member(s1_);
+    uint _e1 = fetch_arg_ptr_member(s1_);
+    assign_to_arg_ptr_array_element(a1_);
+    uint _e3 = fetch_arg_ptr_array_element(a1_);
+    return;
+}
+
 bool index_ptr(bool value) {
     bool a_1[1] = bool[1](false);
     a_1 = bool[1](value);

naga/tests/out/glsl/wgsl-access.foo_vert.Vertex.glsl

@@ -135,6 +135,14 @@ void assign_array_through_ptr_fn(inout vec4 foo_2[2]) {
     return;
 }
 
+void assign_through_ptr() {
+    uint val = 33u;
+    vec4 arr[2] = vec4[2](vec4(6.0), vec4(7.0));
+    assign_through_ptr_fn(val);
+    assign_array_through_ptr_fn(arr);
+    return;
+}
+
 uint fetch_arg_ptr_member(inout AssignToMember p_1) {
     uint _e2 = p_1.x;
     return _e2;
@@ -155,6 +163,16 @@ void assign_to_arg_ptr_array_element(inout uint p_4[4]) {
     return;
 }
 
+void assign_to_ptr_components() {
+    AssignToMember s1_ = AssignToMember(0u);
+    uint a1_[4] = uint[4](0u, 0u, 0u, 0u);
+    assign_to_arg_ptr_member(s1_);
+    uint _e1 = fetch_arg_ptr_member(s1_);
+    assign_to_arg_ptr_array_element(a1_);
+    uint _e3 = fetch_arg_ptr_array_element(a1_);
+    return;
+}
+
 bool index_ptr(bool value) {
     bool a_1[1] = bool[1](false);
     a_1 = bool[1](value);

naga/tests/out/glsl/wgsl-bounds-check-image-restrict.fragment_shader.Fragment.glsl

@@ -10,13 +10,13 @@ uniform sampler3D _group_0_binding_3_fs;
 
 uniform sampler2DMS _group_0_binding_4_fs;
 
-layout(rgba8) writeonly uniform image1D _group_0_binding_8_fs;
+layout(rgba8) writeonly uniform image1D _group_0_binding_5_fs;
 
-layout(rgba8) writeonly uniform image2D _group_0_binding_9_fs;
+layout(rgba8) writeonly uniform image2D _group_0_binding_6_fs;
 
-layout(rgba8) writeonly uniform image2DArray _group_0_binding_10_fs;
+layout(rgba8) writeonly uniform image2DArray _group_0_binding_7_fs;
 
-layout(rgba8) writeonly uniform image3D _group_0_binding_11_fs;
+layout(rgba8) writeonly uniform image3D _group_0_binding_8_fs;
 
 layout(location = 0) out vec4 _fs2p_location0;
 
@@ -56,28 +56,28 @@ vec4 test_textureLoad_multisampled_2d(ivec2 coords_5, int _sample) {
     return _e3;
 }
 
-void test_textureStore_1d(int coords_10, vec4 value) {
+void test_textureStore_1d(int coords_6, vec4 value) {
-    imageStore(_group_0_binding_8_fs, coords_10, value);
+    imageStore(_group_0_binding_5_fs, coords_6, value);
     return;
 }
 
-void test_textureStore_2d(ivec2 coords_11, vec4 value_1) {
+void test_textureStore_2d(ivec2 coords_7, vec4 value_1) {
-    imageStore(_group_0_binding_9_fs, coords_11, value_1);
+    imageStore(_group_0_binding_6_fs, coords_7, value_1);
     return;
 }
 
-void test_textureStore_2d_array_u(ivec2 coords_12, uint array_index, vec4 value_2) {
+void test_textureStore_2d_array_u(ivec2 coords_8, uint array_index, vec4 value_2) {
-    imageStore(_group_0_binding_10_fs, ivec3(coords_12, array_index), value_2);
+    imageStore(_group_0_binding_7_fs, ivec3(coords_8, array_index), value_2);
     return;
 }
 
-void test_textureStore_2d_array_s(ivec2 coords_13, int array_index_1, vec4 value_3) {
+void test_textureStore_2d_array_s(ivec2 coords_9, int array_index_1, vec4 value_3) {
-    imageStore(_group_0_binding_10_fs, ivec3(coords_13, array_index_1), value_3);
+    imageStore(_group_0_binding_7_fs, ivec3(coords_9, array_index_1), value_3);
     return;
 }
 
-void test_textureStore_3d(ivec3 coords_14, vec4 value_4) {
+void test_textureStore_3d(ivec3 coords_10, vec4 value_4) {
-    imageStore(_group_0_binding_11_fs, coords_14, value_4);
+    imageStore(_group_0_binding_8_fs, coords_10, value_4);
     return;
 }
 

naga/tests/out/glsl/wgsl-bounds-check-image-rzsw.fragment_shader.Fragment.glsl

@@ -10,13 +10,13 @@ uniform sampler3D _group_0_binding_3_fs;
 
 uniform sampler2DMS _group_0_binding_4_fs;
 
-layout(rgba8) writeonly uniform image1D _group_0_binding_8_fs;
+layout(rgba8) writeonly uniform image1D _group_0_binding_5_fs;
 
-layout(rgba8) writeonly uniform image2D _group_0_binding_9_fs;
+layout(rgba8) writeonly uniform image2D _group_0_binding_6_fs;
 
-layout(rgba8) writeonly uniform image2DArray _group_0_binding_10_fs;
+layout(rgba8) writeonly uniform image2DArray _group_0_binding_7_fs;
 
-layout(rgba8) writeonly uniform image3D _group_0_binding_11_fs;
+layout(rgba8) writeonly uniform image3D _group_0_binding_8_fs;
 
 layout(location = 0) out vec4 _fs2p_location0;
 
@@ -50,28 +50,28 @@ vec4 test_textureLoad_multisampled_2d(ivec2 coords_5, int _sample) {
     return _e3;
 }
 
-void test_textureStore_1d(int coords_10, vec4 value) {
+void test_textureStore_1d(int coords_6, vec4 value) {
-    imageStore(_group_0_binding_8_fs, coords_10, value);
+    imageStore(_group_0_binding_5_fs, coords_6, value);
     return;
 }
 
-void test_textureStore_2d(ivec2 coords_11, vec4 value_1) {
+void test_textureStore_2d(ivec2 coords_7, vec4 value_1) {
-    imageStore(_group_0_binding_9_fs, coords_11, value_1);
+    imageStore(_group_0_binding_6_fs, coords_7, value_1);
     return;
 }
 
-void test_textureStore_2d_array_u(ivec2 coords_12, uint array_index, vec4 value_2) {
+void test_textureStore_2d_array_u(ivec2 coords_8, uint array_index, vec4 value_2) {
-    imageStore(_group_0_binding_10_fs, ivec3(coords_12, array_index), value_2);
+    imageStore(_group_0_binding_7_fs, ivec3(coords_8, array_index), value_2);
     return;
 }
 
-void test_textureStore_2d_array_s(ivec2 coords_13, int array_index_1, vec4 value_3) {
+void test_textureStore_2d_array_s(ivec2 coords_9, int array_index_1, vec4 value_3) {
-    imageStore(_group_0_binding_10_fs, ivec3(coords_13, array_index_1), value_3);
+    imageStore(_group_0_binding_7_fs, ivec3(coords_9, array_index_1), value_3);
     return;
 }
 
-void test_textureStore_3d(ivec3 coords_14, vec4 value_4) {
+void test_textureStore_3d(ivec3 coords_10, vec4 value_4) {
-    imageStore(_group_0_binding_11_fs, coords_14, value_4);
+    imageStore(_group_0_binding_8_fs, coords_10, value_4);
     return;
 }
 

naga/tests/out/glsl/wgsl-break-if.main.Compute.glsl

@@ -75,6 +75,10 @@ void breakIfSeparateVariable() {
 }
 
 void main() {
+    breakIfEmpty();
+    breakIfEmptyBody(false);
+    breakIf(false);
+    breakIfSeparateVariable();
     return;
 }
 

naga/tests/out/glsl/wgsl-const-exprs.main.Compute.glsl

@@ -10,15 +10,15 @@ const int THREE = 3;
 const bool TRUE = true;
 const bool FALSE = false;
 const int FOUR = 4;
+const int TEXTURE_KIND_REGULAR = 0;
+const int TEXTURE_KIND_WARP = 1;
+const int TEXTURE_KIND_SKY = 2;
 const int FOUR_ALIAS = 4;
 const int TEST_CONSTANT_ADDITION = 8;
 const int TEST_CONSTANT_ALIAS_ADDITION = 8;
 const float PI = 3.141;
 const float phi_sun = 6.282;
 const vec4 DIV = vec4(0.44444445, 0.0, 0.0, 0.0);
-const int TEXTURE_KIND_REGULAR = 0;
-const int TEXTURE_KIND_WARP = 1;
-const int TEXTURE_KIND_SKY = 2;
 const vec2 add_vec = vec2(4.0, 5.0);
 const bvec2 compare_vec = bvec2(true, false);
 
@@ -66,16 +66,6 @@ void compose_of_constant() {
     return;
 }
 
-void compose_of_splat() {
-    vec4 x_1 = vec4(2.0, 1.0, 1.0, 1.0);
-    return;
-}
-
-void test_local_const() {
-    float arr[2] = float[2](0.0, 0.0);
-    return;
-}
-
 uint map_texture_kind(int texture_kind) {
     switch(texture_kind) {
         case 0: {
@@ -93,6 +83,16 @@ uint map_texture_kind(int texture_kind) {
     }
 }
 
+void compose_of_splat() {
+    vec4 x_1 = vec4(2.0, 1.0, 1.0, 1.0);
+    return;
+}
+
+void test_local_const() {
+    float arr[2] = float[2](0.0, 0.0);
+    return;
+}
+
 void compose_vector_zero_val_binop() {
     ivec3 a = ivec3(1, 1, 1);
     ivec3 b = ivec3(0, 1, 2);
@@ -140,8 +140,14 @@ void main() {
     non_constant_initializers();
     splat_of_constant();
     compose_of_constant();
+    uint _e1 = map_texture_kind(1);
     compose_of_splat();
     test_local_const();
+    compose_vector_zero_val_binop();
+    relational();
+    packed_dot_product();
+    test_local_const();
+    abstract_access(1u);
     return;
 }
 

naga/tests/out/glsl/wgsl-control-flow.main.Compute.glsl

@@ -6,6 +6,74 @@ precision highp int;
 layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
 
 
+void control_flow() {
+    int pos = 0;
+    memoryBarrierBuffer();
+    barrier();
+    memoryBarrierShared();
+    barrier();
+    memoryBarrierImage();
+    barrier();
+    do {
+        pos = 1;
+    } while(false);
+    int _e3 = pos;
+    switch(_e3) {
+        case 1: {
+            pos = 0;
+            break;
+        }
+        case 2: {
+            pos = 1;
+            break;
+        }
+        case 3:
+        case 4: {
+            pos = 2;
+            break;
+        }
+        case 5: {
+            pos = 3;
+            break;
+        }
+        default:
+        case 6: {
+            pos = 4;
+            break;
+        }
+    }
+    switch(0u) {
+        case 0u: {
+            break;
+        }
+        default: {
+            break;
+        }
+    }
+    int _e10 = pos;
+    switch(_e10) {
+        case 1: {
+            pos = 0;
+            break;
+        }
+        case 2: {
+            pos = 1;
+            return;
+        }
+        case 3: {
+            pos = 2;
+            return;
+        }
+        case 4: {
+            return;
+        }
+        default: {
+            pos = 3;
+            return;
+        }
+    }
+}
+
 void switch_default_break(int i) {
     do {
         break;
@@ -185,71 +253,14 @@ void loop_switch_omit_continue_variable_checks(int x_2, int y_1, int z_1, int w)
 }
 
 void main() {
-    uvec3 global_id = gl_GlobalInvocationID;
+    control_flow();
-    int pos = 0;
+    switch_default_break(1);
-    memoryBarrierBuffer();
+    switch_case_break();
-    barrier();
+    switch_selector_type_conversion();
-    memoryBarrierShared();
+    switch_const_expr_case_selectors();
-    barrier();
+    loop_switch_continue(1);
-    memoryBarrierImage();
+    loop_switch_continue_nesting(1, 2, 3);
-    barrier();
+    loop_switch_omit_continue_variable_checks(1, 2, 3, 4);
-    do {
-        pos = 1;
-    } while(false);
-    int _e4 = pos;
-    switch(_e4) {
-        case 1: {
-            pos = 0;
-            break;
-        }
-        case 2: {
-            pos = 1;
-            break;
-        }
-        case 3:
-        case 4: {
-            pos = 2;
-            break;
-        }
-        case 5: {
-            pos = 3;
-            break;
-        }
-        default:
-        case 6: {
-            pos = 4;
-            break;
-        }
-    }
-    switch(0u) {
-        case 0u: {
-            break;
-        }
-        default: {
-            break;
-        }
-    }
-    int _e11 = pos;
-    switch(_e11) {
-        case 1: {
-            pos = 0;
-            break;
-        }
-        case 2: {
-            pos = 1;
     return;
-        }
-        case 3: {
-            pos = 2;
-            return;
-        }
-        case 4: {
-            return;
-        }
-        default: {
-            pos = 3;
-            return;
-        }
-    }
 }
 

naga/tests/out/glsl/wgsl-operators.main.Compute.glsl

@@ -267,12 +267,14 @@ void main() {
     uvec3 id = gl_WorkGroupID;
     vec4 _e1 = builtins();
     vec4 _e6 = splat(float(id.x), int(id.y));
-    vec3 _e11 = bool_cast(vec3(1.0, 1.0, 1.0));
+    vec2 _e7 = splat_assignment();
+    vec3 _e12 = bool_cast(vec3(1.0, 1.0, 1.0));
     logical();
     arithmetic();
     bit();
     comparison();
     assignment();
+    negation_avoids_prefix_decrement();
     return;
 }
 

naga/tests/out/glsl/wgsl-pointer-function-arg.main.Compute.glsl

@@ -75,6 +75,19 @@ void let_binding(inout int a_1[4], uint i_7) {
 }
 
 void main() {
+    ivec2 vec[4] = ivec2[4](ivec2(0), ivec2(0), ivec2(0), ivec2(0));
+    mat2x2 mat[4] = mat2x2[4](mat2x2(0.0), mat2x2(0.0), mat2x2(0.0), mat2x2(0.0));
+    int arr1d[4] = int[4](0, 0, 0, 0);
+    int arr2d[4][4] = int[4][4](int[4](0, 0, 0, 0), int[4](0, 0, 0, 0), int[4](0, 0, 0, 0), int[4](0, 0, 0, 0));
+    int arr3d[4][4][4] = int[4][4][4](int[4][4](int[4](0, 0, 0, 0), int[4](0, 0, 0, 0), int[4](0, 0, 0, 0), int[4](0, 0, 0, 0)), int[4][4](int[4](0, 0, 0, 0), int[4](0, 0, 0, 0), int[4](0, 0, 0, 0), int[4](0, 0, 0, 0)), int[4][4](int[4](0, 0, 0, 0), int[4](0, 0, 0, 0), int[4](0, 0, 0, 0), int[4](0, 0, 0, 0)), int[4][4](int[4](0, 0, 0, 0), int[4](0, 0, 0, 0), int[4](0, 0, 0, 0), int[4](0, 0, 0, 0)));
+    local_var(1u);
+    mat_vec_ptrs(vec, mat, 1u);
+    argument(arr1d, 1u);
+    argument_nested_x2_(arr2d, 1u, 2u);
+    argument_nested_x3_(arr3d, 1u, 2u);
+    index_from_self(arr1d, 1u);
+    local_var_from_arg(int[4](1, 2, 3, 4), 5u);
+    let_binding(arr1d, 1u);
     return;
 }
 

naga/tests/out/hlsl/wgsl-abstract-types-return.hlsl

@@ -48,5 +48,13 @@ float2 return_vec2f32_const_ai()
 [numthreads(1, 1, 1)]
 void main()
 {
+    const int _e0 = return_i32_ai();
+    const uint _e1 = return_u32_ai();
+    const float _e2 = return_f32_ai();
+    const float _e3 = return_f32_af();
+    const float2 _e4 = return_vec2f32_ai();
+    const float _e5[4] = return_arrf32_ai();
+    const float _e6 = return_const_f32_const_ai();
+    const float2 _e7 = return_vec2f32_const_ai();
     return;
 }

naga/tests/out/hlsl/wgsl-access.hlsl

@@ -247,6 +247,16 @@ void assign_array_through_ptr_fn(inout float4 foo_2[2])
     return;
 }
 
+void assign_through_ptr()
+{
+    uint val = 33u;
+    float4 arr[2] = Constructarray2_float4_((6.0).xxxx, (7.0).xxxx);
+
+    assign_through_ptr_fn(val);
+    assign_array_through_ptr_fn(arr);
+    return;
+}
+
 uint fetch_arg_ptr_member(inout AssignToMember p_1)
 {
     uint _e2 = p_1.x;
@@ -271,6 +281,18 @@ void assign_to_arg_ptr_array_element(inout uint p_4[4])
     return;
 }
 
+void assign_to_ptr_components()
+{
+    AssignToMember s1_ = (AssignToMember)0;
+    uint a1_[4] = (uint[4])0;
+
+    assign_to_arg_ptr_member(s1_);
+    const uint _e1 = fetch_arg_ptr_member(s1_);
+    assign_to_arg_ptr_array_element(a1_);
+    const uint _e3 = fetch_arg_ptr_array_element(a1_);
+    return;
+}
+
 typedef bool ret_Constructarray1_bool_[1];
 ret_Constructarray1_bool_ Constructarray1_bool_(bool arg0) {
     bool ret[1] = { arg0 };
@@ -406,25 +428,13 @@ float4 foo_frag() : SV_Target0
 }
 
 [numthreads(1, 1, 1)]
-void assign_through_ptr()
+void foo_compute()
 {
-    uint val = 33u;
+    assign_through_ptr();
-    float4 arr[2] = Constructarray2_float4_((6.0).xxxx, (7.0).xxxx);
+    assign_to_ptr_components();
-
+    const bool _e1 = index_ptr(true);
-    assign_through_ptr_fn(val);
+    const int _e2 = member_ptr();
-    assign_array_through_ptr_fn(arr);
+    const int _e3 = let_members_of_members();
-    return;
+    const int _e4 = var_members_of_members();
-}
-
-[numthreads(1, 1, 1)]
-void assign_to_ptr_components()
-{
-    AssignToMember s1_ = (AssignToMember)0;
-    uint a1_[4] = (uint[4])0;
-
-    assign_to_arg_ptr_member(s1_);
-    const uint _e1 = fetch_arg_ptr_member(s1_);
-    assign_to_arg_ptr_array_element(a1_);
-    const uint _e3 = fetch_arg_ptr_array_element(a1_);
     return;
 }

naga/tests/out/hlsl/wgsl-access.ron

@@ -13,11 +13,7 @@
     ],
     compute:[
         (
-            entry_point:"assign_through_ptr",
+            entry_point:"foo_compute",
-            target_profile:"cs_5_1",
-        ),
-        (
-            entry_point:"assign_to_ptr_components",
             target_profile:"cs_5_1",
         ),
     ],

naga/tests/out/hlsl/wgsl-break-if.hlsl

@@ -88,5 +88,9 @@ void breakIfSeparateVariable()
 [numthreads(1, 1, 1)]
 void main()
 {
+    breakIfEmpty();
+    breakIfEmptyBody(false);
+    breakIf(false);
+    breakIfSeparateVariable();
     return;
 }

naga/tests/out/hlsl/wgsl-const-exprs.hlsl

@@ -3,15 +3,15 @@ static const int THREE = int(3);
 static const bool TRUE = true;
 static const bool FALSE = false;
 static const int FOUR = int(4);
+static const int TEXTURE_KIND_REGULAR = int(0);
+static const int TEXTURE_KIND_WARP = int(1);
+static const int TEXTURE_KIND_SKY = int(2);
 static const int FOUR_ALIAS = int(4);
 static const int TEST_CONSTANT_ADDITION = int(8);
 static const int TEST_CONSTANT_ALIAS_ADDITION = int(8);
 static const float PI = 3.141;
 static const float phi_sun = 6.282;
 static const float4 DIV = float4(0.44444445, 0.0, 0.0, 0.0);
-static const int TEXTURE_KIND_REGULAR = int(0);
-static const int TEXTURE_KIND_WARP = int(1);
-static const int TEXTURE_KIND_SKY = int(2);
 static const float2 add_vec = float2(4.0, 5.0);
 static const bool2 compare_vec = bool2(true, false);
 
@@ -70,20 +70,6 @@ void compose_of_constant()
     return;
 }
 
-void compose_of_splat()
-{
-    float4 x_1 = float4(2.0, 1.0, 1.0, 1.0);
-
-    return;
-}
-
-void test_local_const()
-{
-    float arr[2] = (float[2])0;
-
-    return;
-}
-
 uint map_texture_kind(int texture_kind)
 {
     switch(texture_kind) {
@@ -102,6 +88,20 @@ uint map_texture_kind(int texture_kind)
     }
 }
 
+void compose_of_splat()
+{
+    float4 x_1 = float4(2.0, 1.0, 1.0, 1.0);
+
+    return;
+}
+
+void test_local_const()
+{
+    float arr[2] = (float[2])0;
+
+    return;
+}
+
 void compose_vector_zero_val_binop()
 {
     int3 a = int3(int(1), int(1), int(1));
@@ -165,7 +165,13 @@ void main()
     non_constant_initializers();
     splat_of_constant();
     compose_of_constant();
+    const uint _e1 = map_texture_kind(int(1));
     compose_of_splat();
     test_local_const();
+    compose_vector_zero_val_binop();
+    relational();
+    packed_dot_product();
+    test_local_const();
+    abstract_access(1u);
     return;
 }

naga/tests/out/hlsl/wgsl-control-flow.hlsl

@@ -1,3 +1,70 @@
+void control_flow()
+{
+    int pos = (int)0;
+
+    DeviceMemoryBarrierWithGroupSync();
+    GroupMemoryBarrierWithGroupSync();
+    DeviceMemoryBarrierWithGroupSync();
+    do {
+        pos = int(1);
+    } while(false);
+    int _e3 = pos;
+    switch(_e3) {
+        case 1: {
+            pos = int(0);
+            break;
+        }
+        case 2: {
+            pos = int(1);
+            break;
+        }
+        case 3:
+        case 4: {
+            pos = int(2);
+            break;
+        }
+        case 5: {
+            pos = int(3);
+            break;
+        }
+        default:
+        case 6: {
+            pos = int(4);
+            break;
+        }
+    }
+    switch(0u) {
+        case 0u: {
+            break;
+        }
+        default: {
+            break;
+        }
+    }
+    int _e10 = pos;
+    switch(_e10) {
+        case 1: {
+            pos = int(0);
+            break;
+        }
+        case 2: {
+            pos = int(1);
+            return;
+        }
+        case 3: {
+            pos = int(2);
+            return;
+        }
+        case 4: {
+            return;
+        }
+        default: {
+            pos = int(3);
+            return;
+        }
+    }
+}
+
 void switch_default_break(int i)
 {
     do {
@@ -231,69 +298,15 @@ void loop_switch_omit_continue_variable_checks(int x_2, int y_1, int z_1, int w)
 }
 
 [numthreads(1, 1, 1)]
-void main(uint3 global_id : SV_DispatchThreadID)
+void main()
 {
-    int pos = (int)0;
+    control_flow();
-
+    switch_default_break(int(1));
-    DeviceMemoryBarrierWithGroupSync();
+    switch_case_break();
-    GroupMemoryBarrierWithGroupSync();
+    switch_selector_type_conversion();
-    DeviceMemoryBarrierWithGroupSync();
+    switch_const_expr_case_selectors();
-    do {
+    loop_switch_continue(int(1));
-        pos = int(1);
+    loop_switch_continue_nesting(int(1), int(2), int(3));
-    } while(false);
+    loop_switch_omit_continue_variable_checks(int(1), int(2), int(3), int(4));
-    int _e4 = pos;
-    switch(_e4) {
-        case 1: {
-            pos = int(0);
-            break;
-        }
-        case 2: {
-            pos = int(1);
-            break;
-        }
-        case 3:
-        case 4: {
-            pos = int(2);
-            break;
-        }
-        case 5: {
-            pos = int(3);
-            break;
-        }
-        default:
-        case 6: {
-            pos = int(4);
-            break;
-        }
-    }
-    switch(0u) {
-        case 0u: {
-            break;
-        }
-        default: {
-            break;
-        }
-    }
-    int _e11 = pos;
-    switch(_e11) {
-        case 1: {
-            pos = int(0);
-            break;
-        }
-        case 2: {
-            pos = int(1);
     return;
-        }
-        case 3: {
-            pos = int(2);
-            return;
-        }
-        case 4: {
-            return;
-        }
-        default: {
-            pos = int(3);
-            return;
-        }
-    }
 }

naga/tests/out/hlsl/wgsl-conversion-float-to-int.hlsl

@@ -258,5 +258,30 @@ uint64_t2 test_f64_to_u64_vec(double2 f_23)
 [numthreads(1, 1, 1)]
 void main()
 {
+    test_const_eval();
+    const int _e1 = test_f16_to_i32_(1.0h);
+    const uint _e3 = test_f16_to_u32_(1.0h);
+    const int64_t _e5 = test_f16_to_i64_(1.0h);
+    const uint64_t _e7 = test_f16_to_u64_(1.0h);
+    const int _e9 = test_f32_to_i32_(1.0);
+    const uint _e11 = test_f32_to_u32_(1.0);
+    const int64_t _e13 = test_f32_to_i64_(1.0);
+    const uint64_t _e15 = test_f32_to_u64_(1.0);
+    const int _e17 = test_f64_to_i32_(1.0L);
+    const uint _e19 = test_f64_to_u32_(1.0L);
+    const int64_t _e21 = test_f64_to_i64_(1.0L);
+    const uint64_t _e23 = test_f64_to_u64_(1.0L);
+    const int2 _e27 = test_f16_to_i32_vec(half2(1.0h, 2.0h));
+    const uint2 _e31 = test_f16_to_u32_vec(half2(1.0h, 2.0h));
+    const int64_t2 _e35 = test_f16_to_i64_vec(half2(1.0h, 2.0h));
+    const uint64_t2 _e39 = test_f16_to_u64_vec(half2(1.0h, 2.0h));
+    const int2 _e43 = test_f32_to_i32_vec(float2(1.0, 2.0));
+    const uint2 _e47 = test_f32_to_u32_vec(float2(1.0, 2.0));
+    const int64_t2 _e51 = test_f32_to_i64_vec(float2(1.0, 2.0));
+    const uint64_t2 _e55 = test_f32_to_u64_vec(float2(1.0, 2.0));
+    const int2 _e59 = test_f64_to_i32_vec(double2(1.0L, 2.0L));
+    const uint2 _e63 = test_f64_to_u32_vec(double2(1.0L, 2.0L));
+    const int64_t2 _e67 = test_f64_to_i64_vec(double2(1.0L, 2.0L));
+    const uint64_t2 _e71 = test_f64_to_u64_vec(double2(1.0L, 2.0L));
     return;
 }

naga/tests/out/hlsl/wgsl-operators.hlsl

@@ -344,11 +344,13 @@ void main(uint3 id : SV_GroupID)
 {
     const float4 _e1 = builtins();
     const float4 _e6 = splat(float(id.x), int(id.y));
-    const float3 _e11 = bool_cast(float3(1.0, 1.0, 1.0));
+    const float2 _e7 = splat_assignment();
+    const float3 _e12 = bool_cast(float3(1.0, 1.0, 1.0));
     logical();
     arithmetic();
     bit();
     comparison();
     assignment();
+    negation_avoids_prefix_decrement();
     return;
 }

naga/tests/out/hlsl/wgsl-pointer-function-arg.hlsl

@@ -89,5 +89,19 @@ void let_binding(inout int a_1[4], uint i_7)
 [numthreads(1, 1, 1)]
 void main()
 {
+    int2 vec[4] = (int2[4])0;
+    float2x2 mat[4] = (float2x2[4])0;
+    int arr1d[4] = (int[4])0;
+    int arr2d[4][4] = (int[4][4])0;
+    int arr3d[4][4][4] = (int[4][4][4])0;
+
+    local_var(1u);
+    mat_vec_ptrs(vec, mat, 1u);
+    argument(arr1d, 1u);
+    argument_nested_x2_(arr2d, 1u, 2u);
+    argument_nested_x3_(arr3d, 1u, 2u);
+    index_from_self(arr1d, 1u);
+    local_var_from_arg(Constructarray4_int_(int(1), int(2), int(3), int(4)), 5u);
+    let_binding(arr1d, 1u);
     return;
 }

naga/tests/out/ir/wgsl-access.compact.ron

@@ -1761,6 +1761,82 @@
             ],
             diagnostic_filter_leaf: None,
         ),
+        (
+            name: Some("assign_through_ptr"),
+            arguments: [],
+            result: None,
+            local_variables: [
+                (
+                    name: Some("val"),
+                    ty: 0,
+                    init: Some(0),
+                ),
+                (
+                    name: Some("arr"),
+                    ty: 27,
+                    init: Some(6),
+                ),
+            ],
+            expressions: [
+                Literal(U32(33)),
+                LocalVariable(0),
+                Literal(F32(6.0)),
+                Splat(
+                    size: Quad,
+                    value: 2,
+                ),
+                Literal(F32(7.0)),
+                Splat(
+                    size: Quad,
+                    value: 4,
+                ),
+                Compose(
+                    ty: 27,
+                    components: [
+                        3,
+                        5,
+                    ],
+                ),
+                LocalVariable(1),
+            ],
+            named_expressions: {},
+            body: [
+                Call(
+                    function: 4,
+                    arguments: [
+                        1,
+                    ],
+                    result: None,
+                ),
+                Emit((
+                    start: 0,
+                    end: 0,
+                )),
+                Emit((
+                    start: 0,
+                    end: 0,
+                )),
+                Emit((
+                    start: 3,
+                    end: 4,
+                )),
+                Emit((
+                    start: 5,
+                    end: 7,
+                )),
+                Call(
+                    function: 5,
+                    arguments: [
+                        7,
+                    ],
+                    result: None,
+                ),
+                Return(
+                    value: None,
+                ),
+            ],
+            diagnostic_filter_leaf: None,
+        ),
         (
             name: Some("fetch_arg_ptr_member"),
             arguments: [
@@ -1911,6 +1987,64 @@
             ],
             diagnostic_filter_leaf: None,
         ),
+        (
+            name: Some("assign_to_ptr_components"),
+            arguments: [],
+            result: None,
+            local_variables: [
+                (
+                    name: Some("s1"),
+                    ty: 29,
+                    init: None,
+                ),
+                (
+                    name: Some("a1"),
+                    ty: 31,
+                    init: None,
+                ),
+            ],
+            expressions: [
+                LocalVariable(0),
+                CallResult(7),
+                LocalVariable(1),
+                CallResult(9),
+            ],
+            named_expressions: {},
+            body: [
+                Call(
+                    function: 8,
+                    arguments: [
+                        0,
+                    ],
+                    result: None,
+                ),
+                Call(
+                    function: 7,
+                    arguments: [
+                        0,
+                    ],
+                    result: Some(1),
+                ),
+                Call(
+                    function: 10,
+                    arguments: [
+                        2,
+                    ],
+                    result: None,
+                ),
+                Call(
+                    function: 9,
+                    arguments: [
+                        2,
+                    ],
+                    result: Some(3),
+                ),
+                Return(
+                    value: None,
+                ),
+            ],
+            diagnostic_filter_leaf: None,
+        ),
         (
             name: Some("index_ptr"),
             arguments: [
@@ -2705,146 +2839,57 @@
             ),
         ),
         (
-            name: "assign_through_ptr",
+            name: "foo_compute",
             stage: Compute,
             early_depth_test: None,
             workgroup_size: (1, 1, 1),
             workgroup_size_overrides: None,
             function: (
-                name: Some("assign_through_ptr"),
+                name: Some("foo_compute"),
                 arguments: [],
                 result: None,
-                local_variables: [
+                local_variables: [],
-                    (
-                        name: Some("val"),
-                        ty: 0,
-                        init: Some(0),
-                    ),
-                    (
-                        name: Some("arr"),
-                        ty: 27,
-                        init: Some(6),
-                    ),
-                ],
                 expressions: [
-                    Literal(U32(33)),
+                    Literal(Bool(true)),
-                    LocalVariable(0),
+                    CallResult(12),
-                    Literal(F32(6.0)),
+                    CallResult(13),
-                    Splat(
+                    CallResult(14),
-                        size: Quad,
+                    CallResult(15),
-                        value: 2,
-                    ),
-                    Literal(F32(7.0)),
-                    Splat(
-                        size: Quad,
-                        value: 4,
-                    ),
-                    Compose(
-                        ty: 27,
-                        components: [
-                            3,
-                            5,
-                        ],
-                    ),
-                    LocalVariable(1),
                 ],
                 named_expressions: {},
                 body: [
-                    Call(
-                        function: 4,
-                        arguments: [
-                            1,
-                        ],
-                        result: None,
-                    ),
-                    Emit((
-                        start: 0,
-                        end: 0,
-                    )),
-                    Emit((
-                        start: 0,
-                        end: 0,
-                    )),
-                    Emit((
-                        start: 3,
-                        end: 4,
-                    )),
-                    Emit((
-                        start: 5,
-                        end: 7,
-                    )),
-                    Call(
-                        function: 5,
-                        arguments: [
-                            7,
-                        ],
-                        result: None,
-                    ),
-                    Return(
-                        value: None,
-                    ),
-                ],
-                diagnostic_filter_leaf: None,
-            ),
-        ),
-        (
-            name: "assign_to_ptr_components",
-            stage: Compute,
-            early_depth_test: None,
-            workgroup_size: (1, 1, 1),
-            workgroup_size_overrides: None,
-            function: (
-                name: Some("assign_to_ptr_components"),
-                arguments: [],
-                result: None,
-                local_variables: [
-                    (
-                        name: Some("s1"),
-                        ty: 29,
-                        init: None,
-                    ),
-                    (
-                        name: Some("a1"),
-                        ty: 31,
-                        init: None,
-                    ),
-                ],
-                expressions: [
-                    LocalVariable(0),
-                    CallResult(6),
-                    LocalVariable(1),
-                    CallResult(8),
-                ],
-                named_expressions: {},
-                body: [
-                    Call(
-                        function: 7,
-                        arguments: [
-                            0,
-                        ],
-                        result: None,
-                    ),
                     Call(
                         function: 6,
+                        arguments: [],
+                        result: None,
+                    ),
+                    Call(
+                        function: 11,
+                        arguments: [],
+                        result: None,
+                    ),
+                    Call(
+                        function: 12,
                         arguments: [
                             0,
                         ],
                         result: Some(1),
                     ),
                     Call(
-                        function: 9,
+                        function: 13,
-                        arguments: [
+                        arguments: [],
-                            2,
+                        result: Some(2),
-                        ],
-                        result: None,
                     ),
                     Call(
-                        function: 8,
+                        function: 14,
-                        arguments: [
+                        arguments: [],
-                            2,
-                        ],
                         result: Some(3),
                     ),
+                    Call(
+                        function: 15,
+                        arguments: [],
+                        result: Some(4),
+                    ),
                     Return(
                         value: None,
                     ),

naga/tests/out/ir/wgsl-access.ron

@@ -1761,6 +1761,82 @@
             ],
             diagnostic_filter_leaf: None,
         ),
+        (
+            name: Some("assign_through_ptr"),
+            arguments: [],
+            result: None,
+            local_variables: [
+                (
+                    name: Some("val"),
+                    ty: 0,
+                    init: Some(0),
+                ),
+                (
+                    name: Some("arr"),
+                    ty: 27,
+                    init: Some(6),
+                ),
+            ],
+            expressions: [
+                Literal(U32(33)),
+                LocalVariable(0),
+                Literal(F32(6.0)),
+                Splat(
+                    size: Quad,
+                    value: 2,
+                ),
+                Literal(F32(7.0)),
+                Splat(
+                    size: Quad,
+                    value: 4,
+                ),
+                Compose(
+                    ty: 27,
+                    components: [
+                        3,
+                        5,
+                    ],
+                ),
+                LocalVariable(1),
+            ],
+            named_expressions: {},
+            body: [
+                Call(
+                    function: 4,
+                    arguments: [
+                        1,
+                    ],
+                    result: None,
+                ),
+                Emit((
+                    start: 0,
+                    end: 0,
+                )),
+                Emit((
+                    start: 0,
+                    end: 0,
+                )),
+                Emit((
+                    start: 3,
+                    end: 4,
+                )),
+                Emit((
+                    start: 5,
+                    end: 7,
+                )),
+                Call(
+                    function: 5,
+                    arguments: [
+                        7,
+                    ],
+                    result: None,
+                ),
+                Return(
+                    value: None,
+                ),
+            ],
+            diagnostic_filter_leaf: None,
+        ),
         (
             name: Some("fetch_arg_ptr_member"),
             arguments: [
@@ -1911,6 +1987,64 @@
             ],
             diagnostic_filter_leaf: None,
         ),
+        (
+            name: Some("assign_to_ptr_components"),
+            arguments: [],
+            result: None,
+            local_variables: [
+                (
+                    name: Some("s1"),
+                    ty: 29,
+                    init: None,
+                ),
+                (
+                    name: Some("a1"),
+                    ty: 31,
+                    init: None,
+                ),
+            ],
+            expressions: [
+                LocalVariable(0),
+                CallResult(7),
+                LocalVariable(1),
+                CallResult(9),
+            ],
+            named_expressions: {},
+            body: [
+                Call(
+                    function: 8,
+                    arguments: [
+                        0,
+                    ],
+                    result: None,
+                ),
+                Call(
+                    function: 7,
+                    arguments: [
+                        0,
+                    ],
+                    result: Some(1),
+                ),
+                Call(
+                    function: 10,
+                    arguments: [
+                        2,
+                    ],
+                    result: None,
+                ),
+                Call(
+                    function: 9,
+                    arguments: [
+                        2,
+                    ],
+                    result: Some(3),
+                ),
+                Return(
+                    value: None,
+                ),
+            ],
+            diagnostic_filter_leaf: None,
+        ),
         (
             name: Some("index_ptr"),
             arguments: [
@@ -2705,146 +2839,57 @@
             ),
         ),
         (
-            name: "assign_through_ptr",
+            name: "foo_compute",
             stage: Compute,
             early_depth_test: None,
             workgroup_size: (1, 1, 1),
             workgroup_size_overrides: None,
             function: (
-                name: Some("assign_through_ptr"),
+                name: Some("foo_compute"),
                 arguments: [],
                 result: None,
-                local_variables: [
+                local_variables: [],
-                    (
-                        name: Some("val"),
-                        ty: 0,
-                        init: Some(0),
-                    ),
-                    (
-                        name: Some("arr"),
-                        ty: 27,
-                        init: Some(6),
-                    ),
-                ],
                 expressions: [
-                    Literal(U32(33)),
+                    Literal(Bool(true)),
-                    LocalVariable(0),
+                    CallResult(12),
-                    Literal(F32(6.0)),
+                    CallResult(13),
-                    Splat(
+                    CallResult(14),
-                        size: Quad,
+                    CallResult(15),
-                        value: 2,
-                    ),
-                    Literal(F32(7.0)),
-                    Splat(
-                        size: Quad,
-                        value: 4,
-                    ),
-                    Compose(
-                        ty: 27,
-                        components: [
-                            3,
-                            5,
-                        ],
-                    ),
-                    LocalVariable(1),
                 ],
                 named_expressions: {},
                 body: [
-                    Call(
-                        function: 4,
-                        arguments: [
-                            1,
-                        ],
-                        result: None,
-                    ),
-                    Emit((
-                        start: 0,
-                        end: 0,
-                    )),
-                    Emit((
-                        start: 0,
-                        end: 0,
-                    )),
-                    Emit((
-                        start: 3,
-                        end: 4,
-                    )),
-                    Emit((
-                        start: 5,
-                        end: 7,
-                    )),
-                    Call(
-                        function: 5,
-                        arguments: [
-                            7,
-                        ],
-                        result: None,
-                    ),
-                    Return(
-                        value: None,
-                    ),
-                ],
-                diagnostic_filter_leaf: None,
-            ),
-        ),
-        (
-            name: "assign_to_ptr_components",
-            stage: Compute,
-            early_depth_test: None,
-            workgroup_size: (1, 1, 1),
-            workgroup_size_overrides: None,
-            function: (
-                name: Some("assign_to_ptr_components"),
-                arguments: [],
-                result: None,
-                local_variables: [
-                    (
-                        name: Some("s1"),
-                        ty: 29,
-                        init: None,
-                    ),
-                    (
-                        name: Some("a1"),
-                        ty: 31,
-                        init: None,
-                    ),
-                ],
-                expressions: [
-                    LocalVariable(0),
-                    CallResult(6),
-                    LocalVariable(1),
-                    CallResult(8),
-                ],
-                named_expressions: {},
-                body: [
-                    Call(
-                        function: 7,
-                        arguments: [
-                            0,
-                        ],
-                        result: None,
-                    ),
                     Call(
                         function: 6,
+                        arguments: [],
+                        result: None,
+                    ),
+                    Call(
+                        function: 11,
+                        arguments: [],
+                        result: None,
+                    ),
+                    Call(
+                        function: 12,
                         arguments: [
                             0,
                         ],
                         result: Some(1),
                     ),
                     Call(
-                        function: 9,
+                        function: 13,
-                        arguments: [
+                        arguments: [],
-                            2,
+                        result: Some(2),
-                        ],
-                        result: None,
                     ),
                     Call(
-                        function: 8,
+                        function: 14,
-                        arguments: [
+                        arguments: [],
-                            2,
-                        ],
                         result: Some(3),
                     ),
+                    Call(
+                        function: 15,
+                        arguments: [],
+                        result: Some(4),
+                    ),
                     Return(
                         value: None,
                     ),

naga/tests/out/ir/wgsl-const_assert.compact.ron

@@ -10,8 +10,15 @@
     overrides: [],
     global_variables: [],
     global_expressions: [],
-    functions: [
+    functions: [],
+    entry_points: [
         (
+            name: "foo",
+            stage: Compute,
+            early_depth_test: None,
+            workgroup_size: (1, 1, 1),
+            workgroup_size_overrides: None,
+            function: (
                 name: Some("foo"),
                 arguments: [],
                 result: None,
@@ -25,8 +32,8 @@
                 ],
                 diagnostic_filter_leaf: None,
             ),
+        ),
     ],
-    entry_points: [],
     diagnostic_filters: [],
     diagnostic_filter_leaf: None,
 )

naga/tests/out/ir/wgsl-const_assert.ron

@@ -10,8 +10,15 @@
     overrides: [],
     global_variables: [],
     global_expressions: [],
-    functions: [
+    functions: [],
+    entry_points: [
         (
+            name: "foo",
+            stage: Compute,
+            early_depth_test: None,
+            workgroup_size: (1, 1, 1),
+            workgroup_size_overrides: None,
+            function: (
                 name: Some("foo"),
                 arguments: [],
                 result: None,
@@ -25,8 +32,8 @@
                 ],
                 diagnostic_filter_leaf: None,
             ),
+        ),
     ],
-    entry_points: [],
     diagnostic_filters: [],
     diagnostic_filter_leaf: None,
 )

naga/tests/out/ir/wgsl-diagnostic-filter.compact.ron

@@ -40,7 +40,39 @@
             diagnostic_filter_leaf: Some(1),
         ),
     ],
-    entry_points: [],
+    entry_points: [
+        (
+            name: "main",
+            stage: Compute,
+            early_depth_test: None,
+            workgroup_size: (1, 1, 1),
+            workgroup_size_overrides: None,
+            function: (
+                name: Some("main"),
+                arguments: [],
+                result: None,
+                local_variables: [],
+                expressions: [],
+                named_expressions: {},
+                body: [
+                    Call(
+                        function: 0,
+                        arguments: [],
+                        result: None,
+                    ),
+                    Call(
+                        function: 1,
+                        arguments: [],
+                        result: None,
+                    ),
+                    Return(
+                        value: None,
+                    ),
+                ],
+                diagnostic_filter_leaf: Some(0),
+            ),
+        ),
+    ],
     diagnostic_filters: [
         (
             inner: (

naga/tests/out/ir/wgsl-diagnostic-filter.ron

@@ -40,7 +40,39 @@
             diagnostic_filter_leaf: Some(1),
         ),
     ],
-    entry_points: [],
+    entry_points: [
+        (
+            name: "main",
+            stage: Compute,
+            early_depth_test: None,
+            workgroup_size: (1, 1, 1),
+            workgroup_size_overrides: None,
+            function: (
+                name: Some("main"),
+                arguments: [],
+                result: None,
+                local_variables: [],
+                expressions: [],
+                named_expressions: {},
+                body: [
+                    Call(
+                        function: 0,
+                        arguments: [],
+                        result: None,
+                    ),
+                    Call(
+                        function: 1,
+                        arguments: [],
+                        result: None,
+                    ),
+                    Return(
+                        value: None,
+                    ),
+                ],
+                diagnostic_filter_leaf: Some(0),
+            ),
+        ),
+    ],
     diagnostic_filters: [
         (
             inner: (

naga/tests/out/ir/wgsl-index-by-value.compact.ron

@@ -294,28 +294,18 @@
             ],
             diagnostic_filter_leaf: None,
         ),
-    ],
-    entry_points: [
         (
-            name: "index_let_array_1d",
-            stage: Vertex,
-            early_depth_test: None,
-            workgroup_size: (0, 0, 0),
-            workgroup_size_overrides: None,
-            function: (
             name: Some("index_let_array_1d"),
             arguments: [
                 (
                     name: Some("vi"),
                     ty: 7,
-                        binding: Some(BuiltIn(VertexIndex)),
+                    binding: None,
                 ),
             ],
             result: Some((
                 ty: 8,
-                    binding: Some(BuiltIn(Position(
+                binding: None,
-                        invariant: false,
-                    ))),
             )),
             local_variables: [],
             expressions: [
@@ -373,6 +363,106 @@
             ],
             diagnostic_filter_leaf: None,
         ),
+    ],
+    entry_points: [
+        (
+            name: "main",
+            stage: Vertex,
+            early_depth_test: None,
+            workgroup_size: (0, 0, 0),
+            workgroup_size_overrides: None,
+            function: (
+                name: Some("main"),
+                arguments: [
+                    (
+                        name: Some("vi"),
+                        ty: 7,
+                        binding: Some(BuiltIn(VertexIndex)),
+                    ),
+                ],
+                result: Some((
+                    ty: 8,
+                    binding: Some(BuiltIn(Position(
+                        invariant: false,
+                    ))),
+                )),
+                local_variables: [],
+                expressions: [
+                    FunctionArgument(0),
+                    Literal(I32(1)),
+                    Literal(I32(2)),
+                    Literal(I32(3)),
+                    Literal(I32(4)),
+                    Literal(I32(5)),
+                    Compose(
+                        ty: 1,
+                        components: [
+                            1,
+                            2,
+                            3,
+                            4,
+                            5,
+                        ],
+                    ),
+                    Literal(I32(6)),
+                    CallResult(0),
+                    Literal(I32(1)),
+                    Literal(I32(2)),
+                    CallResult(1),
+                    Literal(I32(1)),
+                    Literal(I32(2)),
+                    CallResult(2),
+                    CallResult(3),
+                ],
+                named_expressions: {
+                    0: "vi",
+                },
+                body: [
+                    Emit((
+                        start: 0,
+                        end: 0,
+                    )),
+                    Emit((
+                        start: 6,
+                        end: 7,
+                    )),
+                    Call(
+                        function: 0,
+                        arguments: [
+                            6,
+                            7,
+                        ],
+                        result: Some(8),
+                    ),
+                    Call(
+                        function: 1,
+                        arguments: [
+                            9,
+                            10,
+                        ],
+                        result: Some(11),
+                    ),
+                    Call(
+                        function: 2,
+                        arguments: [
+                            12,
+                            13,
+                        ],
+                        result: Some(14),
+                    ),
+                    Call(
+                        function: 3,
+                        arguments: [
+                            0,
+                        ],
+                        result: Some(15),
+                    ),
+                    Return(
+                        value: Some(15),
+                    ),
+                ],
+                diagnostic_filter_leaf: None,
+            ),
         ),
     ],
     diagnostic_filters: [],

naga/tests/out/ir/wgsl-index-by-value.ron

@@ -294,28 +294,18 @@
             ],
             diagnostic_filter_leaf: None,
         ),
-    ],
-    entry_points: [
         (
-            name: "index_let_array_1d",
-            stage: Vertex,
-            early_depth_test: None,
-            workgroup_size: (0, 0, 0),
-            workgroup_size_overrides: None,
-            function: (
             name: Some("index_let_array_1d"),
             arguments: [
                 (
                     name: Some("vi"),
                     ty: 7,
-                        binding: Some(BuiltIn(VertexIndex)),
+                    binding: None,
                 ),
             ],
             result: Some((
                 ty: 8,
-                    binding: Some(BuiltIn(Position(
+                binding: None,
-                        invariant: false,
-                    ))),
             )),
             local_variables: [],
             expressions: [
@@ -373,6 +363,106 @@
             ],
             diagnostic_filter_leaf: None,
         ),
+    ],
+    entry_points: [
+        (
+            name: "main",
+            stage: Vertex,
+            early_depth_test: None,
+            workgroup_size: (0, 0, 0),
+            workgroup_size_overrides: None,
+            function: (
+                name: Some("main"),
+                arguments: [
+                    (
+                        name: Some("vi"),
+                        ty: 7,
+                        binding: Some(BuiltIn(VertexIndex)),
+                    ),
+                ],
+                result: Some((
+                    ty: 8,
+                    binding: Some(BuiltIn(Position(
+                        invariant: false,
+                    ))),
+                )),
+                local_variables: [],
+                expressions: [
+                    FunctionArgument(0),
+                    Literal(I32(1)),
+                    Literal(I32(2)),
+                    Literal(I32(3)),
+                    Literal(I32(4)),
+                    Literal(I32(5)),
+                    Compose(
+                        ty: 1,
+                        components: [
+                            1,
+                            2,
+                            3,
+                            4,
+                            5,
+                        ],
+                    ),
+                    Literal(I32(6)),
+                    CallResult(0),
+                    Literal(I32(1)),
+                    Literal(I32(2)),
+                    CallResult(1),
+                    Literal(I32(1)),
+                    Literal(I32(2)),
+                    CallResult(2),
+                    CallResult(3),
+                ],
+                named_expressions: {
+                    0: "vi",
+                },
+                body: [
+                    Emit((
+                        start: 0,
+                        end: 0,
+                    )),
+                    Emit((
+                        start: 6,
+                        end: 7,
+                    )),
+                    Call(
+                        function: 0,
+                        arguments: [
+                            6,
+                            7,
+                        ],
+                        result: Some(8),
+                    ),
+                    Call(
+                        function: 1,
+                        arguments: [
+                            9,
+                            10,
+                        ],
+                        result: Some(11),
+                    ),
+                    Call(
+                        function: 2,
+                        arguments: [
+                            12,
+                            13,
+                        ],
+                        result: Some(14),
+                    ),
+                    Call(
+                        function: 3,
+                        arguments: [
+                            0,
+                        ],
+                        result: Some(15),
+                    ),
+                    Return(
+                        value: Some(15),
+                    ),
+                ],
+                diagnostic_filter_leaf: None,
+            ),
         ),
     ],
     diagnostic_filters: [],

naga/tests/out/ir/wgsl-local-const.compact.ron

@@ -72,7 +72,34 @@
             diagnostic_filter_leaf: None,
         ),
     ],
-    entry_points: [],
+    entry_points: [
+        (
+            name: "main",
+            stage: Compute,
+            early_depth_test: None,
+            workgroup_size: (1, 1, 1),
+            workgroup_size_overrides: None,
+            function: (
+                name: Some("main"),
+                arguments: [],
+                result: None,
+                local_variables: [],
+                expressions: [],
+                named_expressions: {},
+                body: [
+                    Call(
+                        function: 0,
+                        arguments: [],
+                        result: None,
+                    ),
+                    Return(
+                        value: None,
+                    ),
+                ],
+                diagnostic_filter_leaf: None,
+            ),
+        ),
+    ],
     diagnostic_filters: [],
     diagnostic_filter_leaf: None,
 )

naga/tests/out/ir/wgsl-local-const.ron

@@ -72,7 +72,34 @@
             diagnostic_filter_leaf: None,
         ),
     ],
-    entry_points: [],
+    entry_points: [
+        (
+            name: "main",
+            stage: Compute,
+            early_depth_test: None,
+            workgroup_size: (1, 1, 1),
+            workgroup_size_overrides: None,
+            function: (
+                name: Some("main"),
+                arguments: [],
+                result: None,
+                local_variables: [],
+                expressions: [],
+                named_expressions: {},
+                body: [
+                    Call(
+                        function: 0,
+                        arguments: [],
+                        result: None,
+                    ),
+                    Return(
+                        value: None,
+                    ),
+                ],
+                diagnostic_filter_leaf: None,
+            ),
+        ),
+    ],
     diagnostic_filters: [],
     diagnostic_filter_leaf: None,
 )

naga/tests/out/ir/wgsl-must-use.compact.ron

@@ -166,9 +166,33 @@
                 arguments: [],
                 result: None,
                 local_variables: [],
-                expressions: [],
+                expressions: [
+                    CallResult(1),
+                    CallResult(2),
+                    CallResult(3),
+                ],
                 named_expressions: {},
                 body: [
+                    Call(
+                        function: 1,
+                        arguments: [],
+                        result: Some(0),
+                    ),
+                    Call(
+                        function: 2,
+                        arguments: [],
+                        result: Some(1),
+                    ),
+                    Call(
+                        function: 3,
+                        arguments: [],
+                        result: Some(2),
+                    ),
+                    Call(
+                        function: 4,
+                        arguments: [],
+                        result: None,
+                    ),
                     Return(
                         value: None,
                     ),

naga/tests/out/ir/wgsl-must-use.ron

@@ -166,9 +166,33 @@
                 arguments: [],
                 result: None,
                 local_variables: [],
-                expressions: [],
+                expressions: [
+                    CallResult(1),
+                    CallResult(2),
+                    CallResult(3),
+                ],
                 named_expressions: {},
                 body: [
+                    Call(
+                        function: 1,
+                        arguments: [],
+                        result: Some(0),
+                    ),
+                    Call(
+                        function: 2,
+                        arguments: [],
+                        result: Some(1),
+                    ),
+                    Call(
+                        function: 3,
+                        arguments: [],
+                        result: Some(2),
+                    ),
+                    Call(
+                        function: 4,
+                        arguments: [],
+                        result: None,
+                    ),
                     Return(
                         value: None,
                     ),

naga/tests/out/msl/wgsl-abstract-types-builtins.msl

@@ -5,7 +5,7 @@
 using metal::uint;
 
 
-void f(
+kernel void f(
 ) {
     int clamp_aiaiai = 1;
     float clamp_aiaiaf = 1.0;

naga/tests/out/msl/wgsl-abstract-types-function-calls.msl

@@ -81,7 +81,7 @@ void func_f_i(
     return;
 }
 
-void main_(
+kernel void main_(
 ) {
     func_f(0.0);
     func_f(0.0);

naga/tests/out/msl/wgsl-abstract-types-let.msl

@@ -129,3 +129,10 @@ void mixed_constant_and_runtime_arguments(
     metal::float2 xvfsf = metal::float2(_e89);
     return;
 }
+
+kernel void main_(
+) {
+    all_constant_arguments();
+    mixed_constant_and_runtime_arguments();
+    return;
+}

naga/tests/out/msl/wgsl-abstract-types-operators.msl

@@ -105,3 +105,17 @@ void wgpu_4435_(
     uint y = a.inner[as_type<int>(as_type<uint>(1) - as_type<uint>(1))];
     return;
 }
+
+kernel void main_(
+  metal::uint3 __local_invocation_id [[thread_position_in_threadgroup]]
+, threadgroup type_3& a
+) {
+    if (metal::all(__local_invocation_id == metal::uint3(0u))) {
+        a = {};
+    }
+    metal::threadgroup_barrier(metal::mem_flags::mem_threadgroup);
+    runtime_values();
+    wgpu_4445_();
+    wgpu_4435_(a);
+    return;
+}

naga/tests/out/msl/wgsl-abstract-types-return.msl

@@ -50,5 +50,13 @@ metal::float2 return_vec2f32_const_ai(
 
 kernel void main_(
 ) {
+    int _e0 = return_i32_ai();
+    uint _e1 = return_u32_ai();
+    float _e2 = return_f32_ai();
+    float _e3 = return_f32_af();
+    metal::float2 _e4 = return_vec2f32_ai();
+    type_4 _e5 = return_arrf32_ai();
+    float _e6 = return_const_f32_const_ai();
+    metal::float2 _e7 = return_vec2f32_const_ai();
     return;
 }

naga/tests/out/msl/wgsl-abstract-types-texture.msl

@@ -35,3 +35,16 @@ void storage(
     st.write(metal::float4(2.0, 3.0, 4.0, 5.0), metal::uint2(metal::int2(0, 1)));
     return;
 }
+
+fragment void main_(
+  metal::texture2d<float, metal::access::sample> t [[user(fake0)]]
+, metal::sampler s [[user(fake0)]]
+, metal::depth2d<float, metal::access::sample> d [[user(fake0)]]
+, metal::sampler c [[user(fake0)]]
+, metal::texture2d<float, metal::access::read_write> st [[user(fake0)]]
+) {
+    color(t, s);
+    depth(s, d, c);
+    storage(st);
+    return;
+}

naga/tests/out/msl/wgsl-abstract-types-var.msl

@@ -253,3 +253,10 @@ void mixed_constant_and_runtime_arguments(
     xvfsf = metal::float2(_e196);
     return;
 }
+
+kernel void main_(
+) {
+    all_constant_arguments();
+    mixed_constant_and_runtime_arguments();
+    return;
+}

naga/tests/out/msl/wgsl-access.msl

@@ -178,6 +178,15 @@ void assign_array_through_ptr_fn(
     return;
 }
 
+void assign_through_ptr(
+) {
+    uint val = 33u;
+    type_22 arr = type_22 {metal::float4(6.0), metal::float4(7.0)};
+    assign_through_ptr_fn(val);
+    assign_array_through_ptr_fn(arr);
+    return;
+}
+
 uint fetch_arg_ptr_member(
     thread AssignToMember& p_1
 ) {
@@ -206,6 +215,17 @@ void assign_to_arg_ptr_array_element(
     return;
 }
 
+void assign_to_ptr_components(
+) {
+    AssignToMember s1_ = {};
+    type_25 a1_ = {};
+    assign_to_arg_ptr_member(s1_);
+    uint _e1 = fetch_arg_ptr_member(s1_);
+    assign_to_arg_ptr_array_element(a1_);
+    uint _e3 = fetch_arg_ptr_array_element(a1_);
+    return;
+}
+
 bool index_ptr(
     bool value
 ) {
@@ -302,23 +322,13 @@ fragment foo_fragOutput foo_frag(
 }
 
 
-kernel void assign_through_ptr(
+kernel void foo_compute(
 ) {
-    uint val = 33u;
+    assign_through_ptr();
-    type_22 arr = type_22 {metal::float4(6.0), metal::float4(7.0)};
+    assign_to_ptr_components();
-    assign_through_ptr_fn(val);
+    bool _e1 = index_ptr(true);
-    assign_array_through_ptr_fn(arr);
+    int _e2 = member_ptr();
-    return;
+    int _e3 = let_members_of_members();
-}
+    int _e4 = var_members_of_members();
-
-
-kernel void assign_to_ptr_components(
-) {
-    AssignToMember s1_ = {};
-    type_25 a1_ = {};
-    assign_to_arg_ptr_member(s1_);
-    uint _e1 = fetch_arg_ptr_member(s1_);
-    assign_to_arg_ptr_array_element(a1_);
-    uint _e3 = fetch_arg_ptr_array_element(a1_);
     return;
 }

naga/tests/out/msl/wgsl-bounds-check-image-restrict-depth.msl

@@ -0,0 +1,62 @@
+// language: metal1.2
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using metal::uint;
+
+
+float test_textureLoad_depth_2d(
+    metal::int2 coords,
+    int level,
+    metal::depth2d<float, metal::access::sample> image_depth_2d
+) {
+    uint clamped_lod_e3 = metal::min(uint(level), image_depth_2d.get_num_mip_levels() - 1);
+    float _e3 = image_depth_2d.read(metal::min(metal::uint2(coords), metal::uint2(image_depth_2d.get_width(clamped_lod_e3), image_depth_2d.get_height(clamped_lod_e3)) - 1), clamped_lod_e3);
+    return _e3;
+}
+
+float test_textureLoad_depth_2d_array_u(
+    metal::int2 coords_1,
+    uint index,
+    int level_1,
+    metal::depth2d_array<float, metal::access::sample> image_depth_2d_array
+) {
+    uint clamped_lod_e4 = metal::min(uint(level_1), image_depth_2d_array.get_num_mip_levels() - 1);
+    float _e4 = image_depth_2d_array.read(metal::min(metal::uint2(coords_1), metal::uint2(image_depth_2d_array.get_width(clamped_lod_e4), image_depth_2d_array.get_height(clamped_lod_e4)) - 1), metal::min(uint(index), image_depth_2d_array.get_array_size() - 1), clamped_lod_e4);
+    return _e4;
+}
+
+float test_textureLoad_depth_2d_array_s(
+    metal::int2 coords_2,
+    int index_1,
+    int level_2,
+    metal::depth2d_array<float, metal::access::sample> image_depth_2d_array
+) {
+    uint clamped_lod_e4 = metal::min(uint(level_2), image_depth_2d_array.get_num_mip_levels() - 1);
+    float _e4 = image_depth_2d_array.read(metal::min(metal::uint2(coords_2), metal::uint2(image_depth_2d_array.get_width(clamped_lod_e4), image_depth_2d_array.get_height(clamped_lod_e4)) - 1), metal::min(uint(index_1), image_depth_2d_array.get_array_size() - 1), clamped_lod_e4);
+    return _e4;
+}
+
+float test_textureLoad_depth_multisampled_2d(
+    metal::int2 coords_3,
+    int _sample,
+    metal::depth2d_ms<float, metal::access::read> image_depth_multisampled_2d
+) {
+    float _e3 = image_depth_multisampled_2d.read(metal::min(metal::uint2(coords_3), metal::uint2(image_depth_multisampled_2d.get_width(), image_depth_multisampled_2d.get_height()) - 1), metal::min(uint(_sample), image_depth_multisampled_2d.get_num_samples() - 1));
+    return _e3;
+}
+
+struct fragment_shaderOutput {
+    metal::float4 member [[color(0)]];
+};
+fragment fragment_shaderOutput fragment_shader(
+  metal::depth2d<float, metal::access::sample> image_depth_2d [[user(fake0)]]
+, metal::depth2d_array<float, metal::access::sample> image_depth_2d_array [[user(fake0)]]
+, metal::depth2d_ms<float, metal::access::read> image_depth_multisampled_2d [[user(fake0)]]
+) {
+    float _e2 = test_textureLoad_depth_2d(metal::int2 {}, 0, image_depth_2d);
+    float _e6 = test_textureLoad_depth_2d_array_u(metal::int2 {}, 0u, 0, image_depth_2d_array);
+    float _e10 = test_textureLoad_depth_2d_array_s(metal::int2 {}, 0, 0, image_depth_2d_array);
+    float _e13 = test_textureLoad_depth_multisampled_2d(metal::int2 {}, 0, image_depth_multisampled_2d);
+    return fragment_shaderOutput { metal::float4(0.0, 0.0, 0.0, 0.0) };
+}

naga/tests/out/msl/wgsl-bounds-check-image-restrict.msl

@@ -65,91 +65,50 @@ metal::float4 test_textureLoad_multisampled_2d(
     return _e3;
 }
 
-float test_textureLoad_depth_2d(
-    metal::int2 coords_6,
-    int level_5,
-    metal::depth2d<float, metal::access::sample> image_depth_2d
-) {
-    uint clamped_lod_e3 = metal::min(uint(level_5), image_depth_2d.get_num_mip_levels() - 1);
-    float _e3 = image_depth_2d.read(metal::min(metal::uint2(coords_6), metal::uint2(image_depth_2d.get_width(clamped_lod_e3), image_depth_2d.get_height(clamped_lod_e3)) - 1), clamped_lod_e3);
-    return _e3;
-}
-
-float test_textureLoad_depth_2d_array_u(
-    metal::int2 coords_7,
-    uint index_2,
-    int level_6,
-    metal::depth2d_array<float, metal::access::sample> image_depth_2d_array
-) {
-    uint clamped_lod_e4 = metal::min(uint(level_6), image_depth_2d_array.get_num_mip_levels() - 1);
-    float _e4 = image_depth_2d_array.read(metal::min(metal::uint2(coords_7), metal::uint2(image_depth_2d_array.get_width(clamped_lod_e4), image_depth_2d_array.get_height(clamped_lod_e4)) - 1), metal::min(uint(index_2), image_depth_2d_array.get_array_size() - 1), clamped_lod_e4);
-    return _e4;
-}
-
-float test_textureLoad_depth_2d_array_s(
-    metal::int2 coords_8,
-    int index_3,
-    int level_7,
-    metal::depth2d_array<float, metal::access::sample> image_depth_2d_array
-) {
-    uint clamped_lod_e4 = metal::min(uint(level_7), image_depth_2d_array.get_num_mip_levels() - 1);
-    float _e4 = image_depth_2d_array.read(metal::min(metal::uint2(coords_8), metal::uint2(image_depth_2d_array.get_width(clamped_lod_e4), image_depth_2d_array.get_height(clamped_lod_e4)) - 1), metal::min(uint(index_3), image_depth_2d_array.get_array_size() - 1), clamped_lod_e4);
-    return _e4;
-}
-
-float test_textureLoad_depth_multisampled_2d(
-    metal::int2 coords_9,
-    int _sample_1,
-    metal::depth2d_ms<float, metal::access::read> image_depth_multisampled_2d
-) {
-    float _e3 = image_depth_multisampled_2d.read(metal::min(metal::uint2(coords_9), metal::uint2(image_depth_multisampled_2d.get_width(), image_depth_multisampled_2d.get_height()) - 1), metal::min(uint(_sample_1), image_depth_multisampled_2d.get_num_samples() - 1));
-    return _e3;
-}
-
 void test_textureStore_1d(
-    int coords_10,
+    int coords_6,
     metal::float4 value,
     metal::texture1d<float, metal::access::write> image_storage_1d
 ) {
-    image_storage_1d.write(value, uint(coords_10));
+    image_storage_1d.write(value, uint(coords_6));
     return;
 }
 
 void test_textureStore_2d(
-    metal::int2 coords_11,
+    metal::int2 coords_7,
     metal::float4 value_1,
     metal::texture2d<float, metal::access::write> image_storage_2d
 ) {
-    image_storage_2d.write(value_1, metal::uint2(coords_11));
+    image_storage_2d.write(value_1, metal::uint2(coords_7));
     return;
 }
 
 void test_textureStore_2d_array_u(
-    metal::int2 coords_12,
+    metal::int2 coords_8,
     uint array_index,
     metal::float4 value_2,
     metal::texture2d_array<float, metal::access::write> image_storage_2d_array
 ) {
-    image_storage_2d_array.write(value_2, metal::uint2(coords_12), array_index);
+    image_storage_2d_array.write(value_2, metal::uint2(coords_8), array_index);
     return;
 }
 
 void test_textureStore_2d_array_s(
-    metal::int2 coords_13,
+    metal::int2 coords_9,
     int array_index_1,
     metal::float4 value_3,
     metal::texture2d_array<float, metal::access::write> image_storage_2d_array
 ) {
-    image_storage_2d_array.write(value_3, metal::uint2(coords_13), array_index_1);
+    image_storage_2d_array.write(value_3, metal::uint2(coords_9), array_index_1);
     return;
 }
 
 void test_textureStore_3d(
-    metal::int3 coords_14,
+    metal::int3 coords_10,
     metal::float4 value_4,
     metal::texture3d<float, metal::access::write> image_storage_3d
 ) {
-    image_storage_3d.write(value_4, metal::uint3(coords_14));
+    image_storage_3d.write(value_4, metal::uint3(coords_10));
     return;
 }
 

naga/tests/out/msl/wgsl-bounds-check-image-rzsw-depth.msl

@@ -0,0 +1,65 @@
+// language: metal1.2
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using metal::uint;
+struct DefaultConstructible {
+    template<typename T>
+    operator T() && {
+        return T {};
+    }
+};
+
+
+float test_textureLoad_depth_2d(
+    metal::int2 coords,
+    int level,
+    metal::depth2d<float, metal::access::sample> image_depth_2d
+) {
+    float _e3 = (uint(level) < image_depth_2d.get_num_mip_levels() && metal::all(metal::uint2(coords) < metal::uint2(image_depth_2d.get_width(level), image_depth_2d.get_height(level))) ? image_depth_2d.read(metal::uint2(coords), level): DefaultConstructible());
+    return _e3;
+}
+
+float test_textureLoad_depth_2d_array_u(
+    metal::int2 coords_1,
+    uint index,
+    int level_1,
+    metal::depth2d_array<float, metal::access::sample> image_depth_2d_array
+) {
+    float _e4 = (uint(level_1) < image_depth_2d_array.get_num_mip_levels() && uint(index) < image_depth_2d_array.get_array_size() && metal::all(metal::uint2(coords_1) < metal::uint2(image_depth_2d_array.get_width(level_1), image_depth_2d_array.get_height(level_1))) ? image_depth_2d_array.read(metal::uint2(coords_1), index, level_1): DefaultConstructible());
+    return _e4;
+}
+
+float test_textureLoad_depth_2d_array_s(
+    metal::int2 coords_2,
+    int index_1,
+    int level_2,
+    metal::depth2d_array<float, metal::access::sample> image_depth_2d_array
+) {
+    float _e4 = (uint(level_2) < image_depth_2d_array.get_num_mip_levels() && uint(index_1) < image_depth_2d_array.get_array_size() && metal::all(metal::uint2(coords_2) < metal::uint2(image_depth_2d_array.get_width(level_2), image_depth_2d_array.get_height(level_2))) ? image_depth_2d_array.read(metal::uint2(coords_2), index_1, level_2): DefaultConstructible());
+    return _e4;
+}
+
+float test_textureLoad_depth_multisampled_2d(
+    metal::int2 coords_3,
+    int _sample,
+    metal::depth2d_ms<float, metal::access::read> image_depth_multisampled_2d
+) {
+    float _e3 = (uint(_sample) < image_depth_multisampled_2d.get_num_samples() && metal::all(metal::uint2(coords_3) < metal::uint2(image_depth_multisampled_2d.get_width(), image_depth_multisampled_2d.get_height())) ? image_depth_multisampled_2d.read(metal::uint2(coords_3), _sample): DefaultConstructible());
+    return _e3;
+}
+
+struct fragment_shaderOutput {
+    metal::float4 member [[color(0)]];
+};
+fragment fragment_shaderOutput fragment_shader(
+  metal::depth2d<float, metal::access::sample> image_depth_2d [[user(fake0)]]
+, metal::depth2d_array<float, metal::access::sample> image_depth_2d_array [[user(fake0)]]
+, metal::depth2d_ms<float, metal::access::read> image_depth_multisampled_2d [[user(fake0)]]
+) {
+    float _e2 = test_textureLoad_depth_2d(metal::int2 {}, 0, image_depth_2d);
+    float _e6 = test_textureLoad_depth_2d_array_u(metal::int2 {}, 0u, 0, image_depth_2d_array);
+    float _e10 = test_textureLoad_depth_2d_array_s(metal::int2 {}, 0, 0, image_depth_2d_array);
+    float _e13 = test_textureLoad_depth_multisampled_2d(metal::int2 {}, 0, image_depth_multisampled_2d);
+    return fragment_shaderOutput { metal::float4(0.0, 0.0, 0.0, 0.0) };
+}

naga/tests/out/msl/wgsl-bounds-check-image-rzsw.msl

@@ -67,88 +67,50 @@ metal::float4 test_textureLoad_multisampled_2d(
     return _e3;
 }
 
-float test_textureLoad_depth_2d(
-    metal::int2 coords_6,
-    int level_5,
-    metal::depth2d<float, metal::access::sample> image_depth_2d
-) {
-    float _e3 = (uint(level_5) < image_depth_2d.get_num_mip_levels() && metal::all(metal::uint2(coords_6) < metal::uint2(image_depth_2d.get_width(level_5), image_depth_2d.get_height(level_5))) ? image_depth_2d.read(metal::uint2(coords_6), level_5): DefaultConstructible());
-    return _e3;
-}
-
-float test_textureLoad_depth_2d_array_u(
-    metal::int2 coords_7,
-    uint index_2,
-    int level_6,
-    metal::depth2d_array<float, metal::access::sample> image_depth_2d_array
-) {
-    float _e4 = (uint(level_6) < image_depth_2d_array.get_num_mip_levels() && uint(index_2) < image_depth_2d_array.get_array_size() && metal::all(metal::uint2(coords_7) < metal::uint2(image_depth_2d_array.get_width(level_6), image_depth_2d_array.get_height(level_6))) ? image_depth_2d_array.read(metal::uint2(coords_7), index_2, level_6): DefaultConstructible());
-    return _e4;
-}
-
-float test_textureLoad_depth_2d_array_s(
-    metal::int2 coords_8,
-    int index_3,
-    int level_7,
-    metal::depth2d_array<float, metal::access::sample> image_depth_2d_array
-) {
-    float _e4 = (uint(level_7) < image_depth_2d_array.get_num_mip_levels() && uint(index_3) < image_depth_2d_array.get_array_size() && metal::all(metal::uint2(coords_8) < metal::uint2(image_depth_2d_array.get_width(level_7), image_depth_2d_array.get_height(level_7))) ? image_depth_2d_array.read(metal::uint2(coords_8), index_3, level_7): DefaultConstructible());
-    return _e4;
-}
-
-float test_textureLoad_depth_multisampled_2d(
-    metal::int2 coords_9,
-    int _sample_1,
-    metal::depth2d_ms<float, metal::access::read> image_depth_multisampled_2d
-) {
-    float _e3 = (uint(_sample_1) < image_depth_multisampled_2d.get_num_samples() && metal::all(metal::uint2(coords_9) < metal::uint2(image_depth_multisampled_2d.get_width(), image_depth_multisampled_2d.get_height())) ? image_depth_multisampled_2d.read(metal::uint2(coords_9), _sample_1): DefaultConstructible());
-    return _e3;
-}
-
 void test_textureStore_1d(
-    int coords_10,
+    int coords_6,
     metal::float4 value,
     metal::texture1d<float, metal::access::write> image_storage_1d
 ) {
-    image_storage_1d.write(value, uint(coords_10));
+    image_storage_1d.write(value, uint(coords_6));
     return;
 }
 
 void test_textureStore_2d(
-    metal::int2 coords_11,
+    metal::int2 coords_7,
     metal::float4 value_1,
     metal::texture2d<float, metal::access::write> image_storage_2d
 ) {
-    image_storage_2d.write(value_1, metal::uint2(coords_11));
+    image_storage_2d.write(value_1, metal::uint2(coords_7));
     return;
 }
 
 void test_textureStore_2d_array_u(
-    metal::int2 coords_12,
+    metal::int2 coords_8,
     uint array_index,
     metal::float4 value_2,
     metal::texture2d_array<float, metal::access::write> image_storage_2d_array
 ) {
-    image_storage_2d_array.write(value_2, metal::uint2(coords_12), array_index);
+    image_storage_2d_array.write(value_2, metal::uint2(coords_8), array_index);
     return;
 }
 
 void test_textureStore_2d_array_s(
-    metal::int2 coords_13,
+    metal::int2 coords_9,
     int array_index_1,
     metal::float4 value_3,
     metal::texture2d_array<float, metal::access::write> image_storage_2d_array
 ) {
-    image_storage_2d_array.write(value_3, metal::uint2(coords_13), array_index_1);
+    image_storage_2d_array.write(value_3, metal::uint2(coords_9), array_index_1);
     return;
 }
 
 void test_textureStore_3d(
-    metal::int3 coords_14,
+    metal::int3 coords_10,
     metal::float4 value_4,
     metal::texture3d<float, metal::access::write> image_storage_3d
 ) {
-    image_storage_3d.write(value_4, metal::uint3(coords_14));
+    image_storage_3d.write(value_4, metal::uint3(coords_10));
     return;
 }
 

naga/tests/out/msl/wgsl-bounds-check-restrict.msl

@@ -184,3 +184,27 @@ void set_dynamic_array_constant_index(
     globals.d[metal::min(unsigned(1000), (_buffer_sizes.size0 - 112 - 4) / 4)] = v_8;
     return;
 }
+
+kernel void main_(
+  device Globals& globals [[user(fake0)]]
+, constant _mslBufferSizes& _buffer_sizes [[user(fake0)]]
+) {
+    float _e1 = index_array(1, globals, _buffer_sizes);
+    float _e3 = index_dynamic_array(1, globals, _buffer_sizes);
+    float _e5 = index_vector(1, globals, _buffer_sizes);
+    float _e12 = index_vector_by_value(metal::float4(2.0, 3.0, 4.0, 5.0), 6);
+    metal::float4 _e14 = index_matrix(1, globals, _buffer_sizes);
+    float _e17 = index_twice(1, 2, globals, _buffer_sizes);
+    float _e19 = index_expensive(1, globals, _buffer_sizes);
+    float _e20 = index_in_bounds(globals, _buffer_sizes);
+    set_array(1, 2.0, globals, _buffer_sizes);
+    set_dynamic_array(1, 2.0, globals, _buffer_sizes);
+    set_vector(1, 2.0, globals, _buffer_sizes);
+    set_matrix(1, metal::float4(2.0, 3.0, 4.0, 5.0), globals, _buffer_sizes);
+    set_index_twice(1, 2, 1.0, globals, _buffer_sizes);
+    set_expensive(1, 1.0, globals, _buffer_sizes);
+    set_in_bounds(1.0, globals, _buffer_sizes);
+    float _e39 = index_dynamic_array_constant_index(globals, _buffer_sizes);
+    set_dynamic_array_constant_index(1.0, globals, _buffer_sizes);
+    return;
+}

naga/tests/out/msl/wgsl-bounds-check-zero-atomic.msl

@@ -91,3 +91,18 @@ uint exchange_atomic_dynamic_sized_array_static_index(
     uint _e4 = uint(1000) < 1 + (_buffer_sizes.size0 - 44 - 4) / 4 ? metal::atomic_exchange_explicit(&globals.c[1000], 1u, metal::memory_order_relaxed) : DefaultConstructible();
     return _e4;
 }
+
+kernel void main_(
+  device Globals& globals [[user(fake0)]]
+, constant _mslBufferSizes& _buffer_sizes [[user(fake0)]]
+) {
+    uint _e0 = fetch_add_atomic(globals, _buffer_sizes);
+    uint _e2 = fetch_add_atomic_static_sized_array(1, globals, _buffer_sizes);
+    uint _e4 = fetch_add_atomic_dynamic_sized_array(1, globals, _buffer_sizes);
+    uint _e5 = exchange_atomic(globals, _buffer_sizes);
+    uint _e7 = exchange_atomic_static_sized_array(1, globals, _buffer_sizes);
+    uint _e9 = exchange_atomic_dynamic_sized_array(1, globals, _buffer_sizes);
+    uint _e10 = fetch_add_atomic_dynamic_sized_array_static_index(globals, _buffer_sizes);
+    uint _e11 = exchange_atomic_dynamic_sized_array_static_index(globals, _buffer_sizes);
+    return;
+}

naga/tests/out/msl/wgsl-bounds-check-zero.msl

@@ -206,3 +206,27 @@ void set_dynamic_array_constant_index(
     }
     return;
 }
+
+kernel void main_(
+  device Globals& globals [[user(fake0)]]
+, constant _mslBufferSizes& _buffer_sizes [[user(fake0)]]
+) {
+    float _e1 = index_array(1, globals, _buffer_sizes);
+    float _e3 = index_dynamic_array(1, globals, _buffer_sizes);
+    float _e5 = index_vector(1, globals, _buffer_sizes);
+    float _e12 = index_vector_by_value(metal::float4(2.0, 3.0, 4.0, 5.0), 6);
+    metal::float4 _e14 = index_matrix(1, globals, _buffer_sizes);
+    float _e17 = index_twice(1, 2, globals, _buffer_sizes);
+    float _e19 = index_expensive(1, globals, _buffer_sizes);
+    float _e20 = index_in_bounds(globals, _buffer_sizes);
+    set_array(1, 2.0, globals, _buffer_sizes);
+    set_dynamic_array(1, 2.0, globals, _buffer_sizes);
+    set_vector(1, 2.0, globals, _buffer_sizes);
+    set_matrix(1, metal::float4(2.0, 3.0, 4.0, 5.0), globals, _buffer_sizes);
+    set_index_twice(1, 2, 1.0, globals, _buffer_sizes);
+    set_expensive(1, 1.0, globals, _buffer_sizes);
+    set_in_bounds(1.0, globals, _buffer_sizes);
+    float _e39 = index_dynamic_array_constant_index(globals, _buffer_sizes);
+    set_dynamic_array_constant_index(1.0, globals, _buffer_sizes);
+    return;
+}

naga/tests/out/msl/wgsl-break-if.msl

@@ -93,5 +93,9 @@ void breakIfSeparateVariable(
 
 kernel void main_(
 ) {
+    breakIfEmpty();
+    breakIfEmptyBody(false);
+    breakIf(false);
+    breakIfSeparateVariable();
     return;
 }

naga/tests/out/msl/wgsl-const-exprs.msl

@@ -7,7 +7,7 @@ using metal::uint;
 struct type_6 {
     float inner[2];
 };
-struct type_10 {
+struct type_9 {
     int inner[9];
 };
 constant uint TWO = 2u;
@@ -15,15 +15,15 @@ constant int THREE = 3;
 constant bool TRUE = true;
 constant bool FALSE = false;
 constant int FOUR = 4;
+constant int TEXTURE_KIND_REGULAR = 0;
+constant int TEXTURE_KIND_WARP = 1;
+constant int TEXTURE_KIND_SKY = 2;
 constant int FOUR_ALIAS = 4;
 constant int TEST_CONSTANT_ADDITION = 8;
 constant int TEST_CONSTANT_ALIAS_ADDITION = 8;
 constant float PI = 3.141;
 constant float phi_sun = 6.282;
 constant metal::float4 DIV = metal::float4(0.44444445, 0.0, 0.0, 0.0);
-constant int TEXTURE_KIND_REGULAR = 0;
-constant int TEXTURE_KIND_WARP = 1;
-constant int TEXTURE_KIND_SKY = 2;
 constant metal::float2 add_vec = metal::float2(4.0, 5.0);
 constant metal::bool2 compare_vec = metal::bool2(true, false);
 
@@ -76,18 +76,6 @@ void compose_of_constant(
     return;
 }
 
-void compose_of_splat(
-) {
-    metal::float4 x_1 = metal::float4(2.0, 1.0, 1.0, 1.0);
-    return;
-}
-
-void test_local_const(
-) {
-    type_6 arr = {};
-    return;
-}
-
 uint map_texture_kind(
     int texture_kind
 ) {
@@ -107,6 +95,18 @@ uint map_texture_kind(
     }
 }
 
+void compose_of_splat(
+) {
+    metal::float4 x_1 = metal::float4(2.0, 1.0, 1.0, 1.0);
+    return;
+}
+
+void test_local_const(
+) {
+    type_6 arr = {};
+    return;
+}
+
 void compose_vector_zero_val_binop(
 ) {
     metal::int3 a = metal::int3(1, 1, 1);
@@ -147,7 +147,7 @@ void abstract_access(
     uint b_1 = 1u;
     int c_1 = {};
     int d = {};
-    c_1 = type_10 {1, 2, 3, 4, 5, 6, 7, 8, 9}.inner[i];
+    c_1 = type_9 {1, 2, 3, 4, 5, 6, 7, 8, 9}.inner[i];
     d = metal::int4(1, 2, 3, 4)[i];
     return;
 }
@@ -160,7 +160,13 @@ kernel void main_(
     non_constant_initializers();
     splat_of_constant();
     compose_of_constant();
+    uint _e1 = map_texture_kind(1);
     compose_of_splat();
     test_local_const();
+    compose_vector_zero_val_binop();
+    relational();
+    packed_dot_product();
+    test_local_const();
+    abstract_access(1u);
     return;
 }

naga/tests/out/msl/wgsl-control-flow.msl

@@ -5,6 +5,75 @@
 using metal::uint;
 
 
+void control_flow(
+) {
+    int pos = {};
+    metal::threadgroup_barrier(metal::mem_flags::mem_device);
+    metal::threadgroup_barrier(metal::mem_flags::mem_threadgroup);
+    metal::threadgroup_barrier(metal::mem_flags::mem_texture);
+    switch(1) {
+        default: {
+            pos = 1;
+            break;
+        }
+    }
+    int _e3 = pos;
+    switch(_e3) {
+        case 1: {
+            pos = 0;
+            break;
+        }
+        case 2: {
+            pos = 1;
+            break;
+        }
+        case 3:
+        case 4: {
+            pos = 2;
+            break;
+        }
+        case 5: {
+            pos = 3;
+            break;
+        }
+        default:
+        case 6: {
+            pos = 4;
+            break;
+        }
+    }
+    switch(0u) {
+        case 0u: {
+            break;
+        }
+        default: {
+            break;
+        }
+    }
+    int _e10 = pos;
+    switch(_e10) {
+        case 1: {
+            pos = 0;
+            break;
+        }
+        case 2: {
+            pos = 1;
+            return;
+        }
+        case 3: {
+            pos = 2;
+            return;
+        }
+        case 4: {
+            return;
+        }
+        default: {
+            pos = 3;
+            return;
+        }
+    }
+}
+
 void switch_default_break(
     int i
 ) {
@@ -214,74 +283,15 @@ void loop_switch_omit_continue_variable_checks(
     return;
 }
 
-struct main_Input {
-};
 kernel void main_(
-  metal::uint3 global_id [[thread_position_in_grid]]
 ) {
-    int pos = {};
+    control_flow();
-    metal::threadgroup_barrier(metal::mem_flags::mem_device);
+    switch_default_break(1);
-    metal::threadgroup_barrier(metal::mem_flags::mem_threadgroup);
+    switch_case_break();
-    metal::threadgroup_barrier(metal::mem_flags::mem_texture);
+    switch_selector_type_conversion();
-    switch(1) {
+    switch_const_expr_case_selectors();
-        default: {
+    loop_switch_continue(1);
-            pos = 1;
+    loop_switch_continue_nesting(1, 2, 3);
-            break;
+    loop_switch_omit_continue_variable_checks(1, 2, 3, 4);
-        }
-    }
-    int _e4 = pos;
-    switch(_e4) {
-        case 1: {
-            pos = 0;
-            break;
-        }
-        case 2: {
-            pos = 1;
-            break;
-        }
-        case 3:
-        case 4: {
-            pos = 2;
-            break;
-        }
-        case 5: {
-            pos = 3;
-            break;
-        }
-        default:
-        case 6: {
-            pos = 4;
-            break;
-        }
-    }
-    switch(0u) {
-        case 0u: {
-            break;
-        }
-        default: {
-            break;
-        }
-    }
-    int _e11 = pos;
-    switch(_e11) {
-        case 1: {
-            pos = 0;
-            break;
-        }
-        case 2: {
-            pos = 1;
     return;
-        }
-        case 3: {
-            pos = 2;
-            return;
-        }
-        case 4: {
-            return;
-        }
-        default: {
-            pos = 3;
-            return;
-        }
-    }
 }

naga/tests/out/msl/wgsl-conversion-float-to-int-no-f64.msl

@@ -197,3 +197,25 @@ metal::ulong2 test_f32_to_u64_vec(
 ) {
     return naga_f2u64(f_15);
 }
+
+kernel void main_(
+) {
+    test_const_eval();
+    int _e1 = test_f16_to_i32_(1.0h);
+    uint _e3 = test_f16_to_u32_(1.0h);
+    long _e5 = test_f16_to_i64_(1.0h);
+    ulong _e7 = test_f16_to_u64_(1.0h);
+    int _e9 = test_f32_to_i32_(1.0);
+    uint _e11 = test_f32_to_u32_(1.0);
+    long _e13 = test_f32_to_i64_(1.0);
+    ulong _e15 = test_f32_to_u64_(1.0);
+    metal::int2 _e19 = test_f16_to_i32_vec(metal::half2(1.0h, 2.0h));
+    metal::uint2 _e23 = test_f16_to_u32_vec(metal::half2(1.0h, 2.0h));
+    metal::long2 _e27 = test_f16_to_i64_vec(metal::half2(1.0h, 2.0h));
+    metal::ulong2 _e31 = test_f16_to_u64_vec(metal::half2(1.0h, 2.0h));
+    metal::int2 _e35 = test_f32_to_i32_vec(metal::float2(1.0, 2.0));
+    metal::uint2 _e39 = test_f32_to_u32_vec(metal::float2(1.0, 2.0));
+    metal::long2 _e43 = test_f32_to_i64_vec(metal::float2(1.0, 2.0));
+    metal::ulong2 _e47 = test_f32_to_u64_vec(metal::float2(1.0, 2.0));
+    return;
+}