More tests for non-struct global types.

src/back/glsl/mod.rs

@@ -923,13 +923,17 @@ impl<'a, W: Write> Writer<'a, W> {
             crate::AddressSpace::PushConstant => {
                 self.write_simple_global(handle, global)?;
             }
-            crate::AddressSpace::Uniform | crate::AddressSpace::Handle => {
+            crate::AddressSpace::Uniform => {
                 self.write_interface_block(handle, global)?;
             }
             crate::AddressSpace::Storage { .. } => {
                 self.write_interface_block(handle, global)?;
             }
+            // A global variable in the `Function` address space is a
+            // contradiction in terms.
             crate::AddressSpace::Function => unreachable!(),
+            // Textures and samplers are handled directly in `Writer::write`.
+            crate::AddressSpace::Handle => unreachable!(),
         }
 
         Ok(())

tests/in/globals.wgsl

@@ -19,6 +19,9 @@ var<storage> dummy: array<vec2<f32>>;
 @group(0) @binding(3)
 var<uniform> float_vecs: array<vec4<f32>, 20>;
 
+@group(0) @binding(4)
+var<uniform> global_vec: vec4<f32>;
+
 fn test_msl_packed_vec3_as_arg(arg: vec3<f32>) {}
 
 fn test_msl_packed_vec3() {
@@ -50,6 +53,9 @@ fn test_msl_packed_vec3() {
 fn main() {
     test_msl_packed_vec3();
 
+    wg[6] = global_vec.x;
+    wg[5] = dummy[1].y;
+    wg[4] = float_vecs[0].w;
     wg[3] = alignment.v1;
     wg[2] = alignment.v3.x;
     alignment.v1 = 4.0;

tests/out/glsl/globals.main.Compute.glsl

@@ -17,6 +17,10 @@ layout(std430) buffer Foo_block_0Compute { Foo _group_0_binding_1_cs; };
 
 layout(std430) readonly buffer type_6_block_1Compute { vec2 _group_0_binding_2_cs[]; };
 
+uniform type_8_block_2Compute { vec4 _group_0_binding_3_cs[20]; };
+
+uniform type_7_block_3Compute { vec4 _group_0_binding_4_cs; };
+
 
 void test_msl_packed_vec3_as_arg(vec3 arg) {
     return;
@@ -27,8 +31,8 @@ void test_msl_packed_vec3_() {
     _group_0_binding_1_cs.v3_ = vec3(1.0);
     _group_0_binding_1_cs.v3_.x = 1.0;
     _group_0_binding_1_cs.v3_.x = 2.0;
-    int _e19 = idx;
-    _group_0_binding_1_cs.v3_[_e19] = 3.0;
+    int _e20 = idx;
+    _group_0_binding_1_cs.v3_[_e20] = 3.0;
     Foo data = _group_0_binding_1_cs;
     vec3 unnamed = data.v3_;
     vec2 unnamed_1 = data.v3_.zx;
@@ -43,10 +47,16 @@ void main() {
     float Foo_1 = 1.0;
     bool at = true;
     test_msl_packed_vec3_();
-    float _e9 = _group_0_binding_1_cs.v1_;
-    wg[3] = _e9;
-    float _e14 = _group_0_binding_1_cs.v3_.x;
-    wg[2] = _e14;
+    float _e10 = _group_0_binding_4_cs.x;
+    wg[6] = _e10;
+    float _e16 = _group_0_binding_2_cs[1].y;
+    wg[5] = _e16;
+    float _e22 = _group_0_binding_3_cs[0].w;
+    wg[4] = _e22;
+    float _e26 = _group_0_binding_1_cs.v1_;
+    wg[3] = _e26;
+    float _e31 = _group_0_binding_1_cs.v3_.x;
+    wg[2] = _e31;
     _group_0_binding_1_cs.v1_ = 4.0;
     wg[1] = float(uint(_group_0_binding_2_cs.length()));
     at_1 = 2u;

tests/out/hlsl/globals.hlsl

@@ -10,6 +10,7 @@ groupshared uint at_1;
 RWByteAddressBuffer alignment : register(u1);
 ByteAddressBuffer dummy : register(t2);
 cbuffer float_vecs : register(b3) { float4 float_vecs[20]; }
+cbuffer global_vec : register(b4) { float4 global_vec; }
 
 void test_msl_packed_vec3_as_arg(float3 arg)
 {
@@ -23,8 +24,8 @@ void test_msl_packed_vec3_()
     alignment.Store3(0, asuint(float3(1.0.xxx)));
     alignment.Store(0+0, asuint(1.0));
     alignment.Store(0+0, asuint(2.0));
-    int _expr19 = idx;
-    alignment.Store(_expr19*4+0, asuint(3.0));
+    int _expr20 = idx;
+    alignment.Store(_expr20*4+0, asuint(3.0));
     Foo data = {asfloat(alignment.Load3(0)), asfloat(alignment.Load(12))};
     float3 unnamed = data.v3_;
     float2 unnamed_1 = data.v3_.zx;
@@ -49,10 +50,16 @@ void main()
     bool at = true;
 
     test_msl_packed_vec3_();
-    float _expr9 = asfloat(alignment.Load(12));
-    wg[3] = _expr9;
-    float _expr14 = asfloat(alignment.Load(0+0));
-    wg[2] = _expr14;
+    float _expr10 = global_vec.x;
+    wg[6] = _expr10;
+    float _expr16 = asfloat(dummy.Load(4+8));
+    wg[5] = _expr16;
+    float _expr22 = float_vecs[0].w;
+    wg[4] = _expr22;
+    float _expr26 = asfloat(alignment.Load(12));
+    wg[3] = _expr26;
+    float _expr31 = asfloat(alignment.Load(0+0));
+    wg[2] = _expr31;
     alignment.Store(12, asuint(4.0));
     wg[1] = float(((NagaBufferLength(dummy) - 0) / 8));
     at_1 = 2u;

tests/out/msl/globals.msl

@@ -36,8 +36,8 @@ void test_msl_packed_vec3_(
     alignment.v3_ = metal::float3(1.0);
     alignment.v3_[0] = 1.0;
     alignment.v3_[0] = 2.0;
-    int _e19 = idx;
-    alignment.v3_[_e19] = 3.0;
+    int _e20 = idx;
+    alignment.v3_[_e20] = 3.0;
     Foo data = alignment;
     metal::float3 unnamed = data.v3_;
     metal::float2 unnamed_1 = metal::float3(data.v3_).zx;
@@ -53,15 +53,23 @@ kernel void main_(
 , threadgroup metal::atomic_uint& at_1
 , device Foo& alignment [[user(fake0)]]
 , device type_6 const& dummy [[user(fake0)]]
+, constant type_8& float_vecs [[user(fake0)]]
+, constant metal::float4& global_vec [[user(fake0)]]
 , constant _mslBufferSizes& _buffer_sizes [[user(fake0)]]
 ) {
     float Foo_1 = 1.0;
     bool at = true;
     test_msl_packed_vec3_(alignment);
-    float _e9 = alignment.v1_;
-    wg.inner[3] = _e9;
-    float _e14 = alignment.v3_[0];
-    wg.inner[2] = _e14;
+    float _e10 = global_vec.x;
+    wg.inner[6] = _e10;
+    float _e16 = dummy[1].y;
+    wg.inner[5] = _e16;
+    float _e22 = float_vecs.inner[0].w;
+    wg.inner[4] = _e22;
+    float _e26 = alignment.v1_;
+    wg.inner[3] = _e26;
+    float _e31 = alignment.v3_[0];
+    wg.inner[2] = _e31;
     alignment.v1_ = 4.0;
     wg.inner[1] = static_cast<float>(1 + (_buffer_sizes.size3 - 0 - 8) / 8);
     metal::atomic_store_explicit(&at_1, 2u, metal::memory_order_relaxed);