Test matrix multiplies in more complex scenarios.

reference/shaders-msl-no-opt/packing/matrix-multiply-unpacked-col-major-2.comp

@@ -0,0 +1,23 @@
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+typedef packed_float3 packed_float3x3[3];
+
+struct SSBO
+{
+    packed_float3x3 m0;
+    packed_float3x3 m1;
+    packed_float3 v0;
+    packed_float3 v1;
+};
+
+kernel void main0(device SSBO& _11 [[buffer(0)]])
+{
+    _11.v0 = (float3x3(float3(_11.m0[0]), float3(_11.m0[1]), float3(_11.m0[2])) * float3x3(float3(_11.m1[0]), float3(_11.m1[1]), float3(_11.m1[2]))) * float3(_11.v1);
+    _11.v0 = float3x3(float3(_11.m0[0]), float3(_11.m0[1]), float3(_11.m0[2])) * (float3x3(float3(_11.m1[0]), float3(_11.m1[1]), float3(_11.m1[2])) * float3(_11.v1));
+    _11.v0 = (float3(_11.v1) * float3x3(float3(_11.m0[0]), float3(_11.m0[1]), float3(_11.m0[2]))) * float3x3(float3(_11.m1[0]), float3(_11.m1[1]), float3(_11.m1[2]));
+    _11.v0 = float3(_11.v1) * (float3x3(float3(_11.m0[0]), float3(_11.m0[1]), float3(_11.m0[2])) * float3x3(float3(_11.m1[0]), float3(_11.m1[1]), float3(_11.m1[2])));
+}
+

reference/shaders-msl-no-opt/packing/matrix-multiply-unpacked-col-major.comp

@@ -0,0 +1,21 @@
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+struct SSBO
+{
+    float2x4 m0;
+    float2x4 m1;
+    float2 v0;
+    float2 v1;
+};
+
+kernel void main0(device SSBO& _11 [[buffer(0)]])
+{
+    _11.v0 = (float2x2(_11.m0[0].xy, _11.m0[1].xy) * float2x2(_11.m1[0].xy, _11.m1[1].xy)) * _11.v1;
+    _11.v0 = float2x2(_11.m0[0].xy, _11.m0[1].xy) * (float2x2(_11.m1[0].xy, _11.m1[1].xy) * _11.v1);
+    _11.v0 = (_11.v1 * float2x2(_11.m0[0].xy, _11.m0[1].xy)) * float2x2(_11.m1[0].xy, _11.m1[1].xy);
+    _11.v0 = _11.v1 * (float2x2(_11.m0[0].xy, _11.m0[1].xy) * float2x2(_11.m1[0].xy, _11.m1[1].xy));
+}
+

reference/shaders-msl-no-opt/packing/matrix-multiply-unpacked-row-major-2.comp

@@ -0,0 +1,23 @@
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+typedef packed_float3 packed_rm_float3x3[3];
+
+struct SSBO
+{
+    packed_rm_float3x3 m0;
+    packed_rm_float3x3 m1;
+    packed_float3 v0;
+    packed_float3 v1;
+};
+
+kernel void main0(device SSBO& _11 [[buffer(0)]])
+{
+    _11.v0 = _11.v1 * (float3x3(float3(_11.m1[0]), float3(_11.m1[1]), float3(_11.m1[2])) * float3x3(float3(_11.m0[0]), float3(_11.m0[1]), float3(_11.m0[2])));
+    _11.v0 = (_11.v1 * float3x3(float3(_11.m1[0]), float3(_11.m1[1]), float3(_11.m1[2]))) * float3x3(float3(_11.m0[0]), float3(_11.m0[1]), float3(_11.m0[2]));
+    _11.v0 = float3x3(float3(_11.m1[0]), float3(_11.m1[1]), float3(_11.m1[2])) * (float3x3(float3(_11.m0[0]), float3(_11.m0[1]), float3(_11.m0[2])) * _11.v1);
+    _11.v0 = (float3x3(float3(_11.m1[0]), float3(_11.m1[1]), float3(_11.m1[2])) * float3x3(float3(_11.m0[0]), float3(_11.m0[1]), float3(_11.m0[2]))) * _11.v1;
+}
+

shaders-msl-no-opt/packing/matrix-multiply-unpacked-col-major-2.comp

@@ -0,0 +1,19 @@
+#version 450
+#extension GL_EXT_scalar_block_layout : require
+layout(local_size_x = 1) in;
+
+layout(scalar, set = 0, binding = 0) buffer SSBO
+{
+	mat3 m0;
+	mat3 m1;
+	vec3 v0;
+	vec3 v1;
+};
+
+void main()
+{
+	v0 = (m0 * m1) * v1;
+	v0 = m0 * (m1 * v1);
+	v0 = (v1 * m0) * m1;
+	v0 = v1 * (m0 * m1);
+}

shaders-msl-no-opt/packing/matrix-multiply-unpacked-col-major.comp

@@ -0,0 +1,18 @@
+#version 450
+layout(local_size_x = 1) in;
+
+layout(std140, set = 0, binding = 0) buffer SSBO
+{
+	mat2 m0;
+	mat2 m1;
+	vec2 v0;
+	vec2 v1;
+};
+
+void main()
+{
+	v0 = (m0 * m1) * v1;
+	v0 = m0 * (m1 * v1);
+	v0 = (v1 * m0) * m1;
+	v0 = v1 * (m0 * m1);
+}

shaders-msl-no-opt/packing/matrix-multiply-unpacked-row-major-2.comp

@@ -0,0 +1,19 @@
+#version 450
+#extension GL_EXT_scalar_block_layout : require
+layout(local_size_x = 1) in;
+
+layout(scalar, row_major, set = 0, binding = 0) buffer SSBO
+{
+	mat3 m0;
+	mat3 m1;
+	vec3 v0;
+	vec3 v1;
+};
+
+void main()
+{
+	v0 = (m0 * m1) * v1;
+	v0 = m0 * (m1 * v1);
+	v0 = (v1 * m0) * m1;
+	v0 = v1 * (m0 * m1);
+}