Test matrix-typed uniform globals.

tests/in/globals.wgsl

@@ -22,6 +22,9 @@ var<uniform> float_vecs: array<vec4<f32>, 20>;
 @group(0) @binding(4)
 var<uniform> global_vec: vec4<f32>;
 
+@group(0) @binding(5)
+var<uniform> global_mat: mat4x4<f32>;
+
 fn test_msl_packed_vec3_as_arg(arg: vec3<f32>) {}
 
 fn test_msl_packed_vec3() {
@@ -53,7 +56,7 @@ fn test_msl_packed_vec3() {
 fn main() {
     test_msl_packed_vec3();
 
-    wg[6] = global_vec.x;
+    wg[6] = (global_mat * global_vec).x;
     wg[5] = dummy[1].y;
     wg[4] = float_vecs[0].w;
     wg[3] = alignment.v1;

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

@@ -21,6 +21,8 @@ uniform type_8_block_2Compute { vec4 _group_0_binding_3_cs[20]; };
 
 uniform type_7_block_3Compute { vec4 _group_0_binding_4_cs; };
 
+uniform type_9_block_4Compute { mat4x4 _group_0_binding_5_cs; };
+
 
 void test_msl_packed_vec3_as_arg(vec3 arg) {
     return;
@@ -31,8 +33,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 _e20 = idx;
-    _group_0_binding_1_cs.v3_[_e20] = 3.0;
+    int _e21 = idx;
+    _group_0_binding_1_cs.v3_[_e21] = 3.0;
     Foo data = _group_0_binding_1_cs;
     vec3 unnamed = data.v3_;
     vec2 unnamed_1 = data.v3_.zx;
@@ -47,16 +49,17 @@ void main() {
     float Foo_1 = 1.0;
     bool at = true;
     test_msl_packed_vec3_();
-    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;
+    mat4x4 _e10 = _group_0_binding_5_cs;
+    vec4 _e11 = _group_0_binding_4_cs;
+    wg[6] = (_e10 * _e11).x;
+    float _e19 = _group_0_binding_2_cs[1].y;
+    wg[5] = _e19;
+    float _e25 = _group_0_binding_3_cs[0].w;
+    wg[4] = _e25;
+    float _e29 = _group_0_binding_1_cs.v1_;
+    wg[3] = _e29;
+    float _e34 = _group_0_binding_1_cs.v3_.x;
+    wg[2] = _e34;
     _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

@@ -11,6 +11,7 @@ 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; }
+cbuffer global_mat : register(b5) { float4x4 global_mat; }
 
 void test_msl_packed_vec3_as_arg(float3 arg)
 {
@@ -24,8 +25,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 _expr20 = idx;
-    alignment.Store(_expr20*4+0, asuint(3.0));
+    int _expr21 = idx;
+    alignment.Store(_expr21*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;
@@ -50,16 +51,17 @@ void main()
     bool at = true;
 
     test_msl_packed_vec3_();
-    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;
+    float4x4 _expr10 = global_mat;
+    float4 _expr11 = global_vec;
+    wg[6] = mul(_expr11, _expr10).x;
+    float _expr19 = asfloat(dummy.Load(4+8));
+    wg[5] = _expr19;
+    float _expr25 = float_vecs[0].w;
+    wg[4] = _expr25;
+    float _expr29 = asfloat(alignment.Load(12));
+    wg[3] = _expr29;
+    float _expr34 = asfloat(alignment.Load(0+0));
+    wg[2] = _expr34;
     alignment.Store(12, asuint(4.0));
     wg[1] = float(((NagaBufferLength(dummy) - 0) / 8));
     at_1 = 2u;

tests/out/msl/globals.msl

@@ -21,7 +21,7 @@ struct type_8 {
     metal::float4 inner[20];
 };
 constant metal::float3 const_type_4_ = {0.0, 0.0, 0.0};
-constant metal::float3x3 const_type_10_ = {const_type_4_, const_type_4_, const_type_4_};
+constant metal::float3x3 const_type_11_ = {const_type_4_, const_type_4_, const_type_4_};
 
 void test_msl_packed_vec3_as_arg(
     metal::float3 arg
@@ -36,14 +36,14 @@ void test_msl_packed_vec3_(
     alignment.v3_ = metal::float3(1.0);
     alignment.v3_[0] = 1.0;
     alignment.v3_[0] = 2.0;
-    int _e20 = idx;
-    alignment.v3_[_e20] = 3.0;
+    int _e21 = idx;
+    alignment.v3_[_e21] = 3.0;
     Foo data = alignment;
     metal::float3 unnamed = data.v3_;
     metal::float2 unnamed_1 = metal::float3(data.v3_).zx;
     test_msl_packed_vec3_as_arg(data.v3_);
-    metal::float3 unnamed_2 = metal::float3(data.v3_) * const_type_10_;
-    metal::float3 unnamed_3 = const_type_10_ * metal::float3(data.v3_);
+    metal::float3 unnamed_2 = metal::float3(data.v3_) * const_type_11_;
+    metal::float3 unnamed_3 = const_type_11_ * metal::float3(data.v3_);
     metal::float3 unnamed_4 = data.v3_ * 2.0;
     metal::float3 unnamed_5 = 2.0 * data.v3_;
 }
@@ -55,21 +55,23 @@ kernel void main_(
 , device type_6 const& dummy [[user(fake0)]]
 , constant type_8& float_vecs [[user(fake0)]]
 , constant metal::float4& global_vec [[user(fake0)]]
+, constant metal::float4x4& global_mat [[user(fake0)]]
 , constant _mslBufferSizes& _buffer_sizes [[user(fake0)]]
 ) {
     float Foo_1 = 1.0;
     bool at = true;
     test_msl_packed_vec3_(alignment);
-    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;
+    metal::float4x4 _e10 = global_mat;
+    metal::float4 _e11 = global_vec;
+    wg.inner[6] = (_e10 * _e11).x;
+    float _e19 = dummy[1].y;
+    wg.inner[5] = _e19;
+    float _e25 = float_vecs.inner[0].w;
+    wg.inner[4] = _e25;
+    float _e29 = alignment.v1_;
+    wg.inner[3] = _e29;
+    float _e34 = alignment.v3_[0];
+    wg.inner[2] = _e34;
     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);