hlsl: support arrays of matrices (gfx-rs#1629)

src/back/hlsl/writer.rs

@@ -715,22 +715,11 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> {
                     stride: _,
                 } => {
                     // HLSL arrays are written as `type name[size]`
-                    let (ty_name, vec_size) = match module.types[base].inner {
-                        // Write scalar type by backend so as not to depend on the front-end implementation
-                        // Name returned from frontend can be generated (type1, float1, etc.)
-                        TypeInner::Scalar { kind, width } => (kind.to_hlsl_str(width)?, None),
-                        // Similarly, write vector types directly.
-                        TypeInner::Vector { size, kind, width } => {
-                            (kind.to_hlsl_str(width)?, Some(size))
-                        }
-                        _ => (self.names[&NameKey::Type(base)].as_str(), None),
-                    };
-
-                    // Write `type` and `name`
-                    write!(self.out, "{}", ty_name)?;
-                    if let Some(s) = vec_size {
-                        write!(self.out, "{}", s as usize)?;
+                    if let TypeInner::Matrix { .. } = module.types[base].inner {
+                        write!(self.out, "row_major ")?;
                     }
+                    self.write_type(module, base)?;
+                    // Write `name`
                     write!(
                         self.out,
                         " {}",

tests/in/access.wgsl

@@ -2,6 +2,7 @@
 
 struct Bar {
 	matrix: mat4x4<f32>;
+	matrix_array: array<mat2x2<f32>, 2>;
 	atom: atomic<i32>;
 	arr: [[stride(8)]] array<vec2<u32>, 2>;
 	data: [[stride(8)]] array<i32>;

tests/out/glsl/access.atomics.Compute.glsl

@@ -7,6 +7,7 @@ layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
 
 layout(std430) buffer Bar_block_0Compute {
     mat4x4 matrix;
+    mat2x2 matrix_array[2];
     int atom;
     uvec2 arr[2];
     int data[];

tests/out/glsl/access.foo.Vertex.glsl

@@ -5,6 +5,7 @@ precision highp int;
 
 layout(std430) buffer Bar_block_0Vertex {
     mat4x4 matrix;
+    mat2x2 matrix_array[2];
     int atom;
     uvec2 arr[2];
     int data[];

tests/out/hlsl/access.hlsl

@@ -22,9 +22,9 @@ float4 foo(uint vi : SV_VertexID) : SV_Position
     float baz = foo_1;
     foo_1 = 1.0;
     float4x4 matrix_ = float4x4(asfloat(bar.Load4(0+0)), asfloat(bar.Load4(0+16)), asfloat(bar.Load4(0+32)), asfloat(bar.Load4(0+48)));
-    uint2 arr[2] = {asuint(bar.Load2(72+0)), asuint(bar.Load2(72+8))};
+    uint2 arr[2] = {asuint(bar.Load2(104+0)), asuint(bar.Load2(104+8))};
     float b = asfloat(bar.Load(0+48+0));
-    int a = asint(bar.Load((((NagaBufferLengthRW(bar) - 88) / 8) - 2u)*8+88));
+    int a = asint(bar.Load((((NagaBufferLengthRW(bar) - 120) / 8) - 2u)*8+120));
     const float _e25 = read_from_private(foo_1);
     bar.Store(8+16+0, asuint(1.0));
     {
@@ -36,10 +36,10 @@ float4 foo(uint vi : SV_VertexID) : SV_Position
     }
     {
         uint2 _value2[2] = { uint2(0u.xx), uint2(1u.xx) };
-        bar.Store2(72+0, asuint(_value2[0]));
-        bar.Store2(72+8, asuint(_value2[1]));
+        bar.Store2(104+0, asuint(_value2[0]));
+        bar.Store2(104+8, asuint(_value2[1]));
     }
-    bar.Store(8+88, asuint(1));
+    bar.Store(8+120, asuint(1));
     {
         int _result[5]={ a, int(b), 3, 4, 5 };
         for(int _i=0; _i<5; ++_i) c[_i] = _result[_i];
@@ -54,23 +54,23 @@ void atomics()
 {
     int tmp = (int)0;
 
-    int value_1 = asint(bar.Load(64));
-    int _e6; bar.InterlockedAdd(64, 5, _e6);
+    int value_1 = asint(bar.Load(96));
+    int _e6; bar.InterlockedAdd(96, 5, _e6);
     tmp = _e6;
-    int _e9; bar.InterlockedAdd(64, -5, _e9);
+    int _e9; bar.InterlockedAdd(96, -5, _e9);
     tmp = _e9;
-    int _e12; bar.InterlockedAnd(64, 5, _e12);
+    int _e12; bar.InterlockedAnd(96, 5, _e12);
     tmp = _e12;
-    int _e15; bar.InterlockedOr(64, 5, _e15);
+    int _e15; bar.InterlockedOr(96, 5, _e15);
     tmp = _e15;
-    int _e18; bar.InterlockedXor(64, 5, _e18);
+    int _e18; bar.InterlockedXor(96, 5, _e18);
     tmp = _e18;
-    int _e21; bar.InterlockedMin(64, 5, _e21);
+    int _e21; bar.InterlockedMin(96, 5, _e21);
     tmp = _e21;
-    int _e24; bar.InterlockedMax(64, 5, _e24);
+    int _e24; bar.InterlockedMax(96, 5, _e24);
     tmp = _e24;
-    int _e27; bar.InterlockedExchange(64, 5, _e27);
+    int _e27; bar.InterlockedExchange(96, 5, _e27);
     tmp = _e27;
-    bar.Store(64, asuint(value_1));
+    bar.Store(96, asuint(value_1));
     return;
 }

tests/out/msl/access.msl

@@ -6,18 +6,22 @@ struct _mslBufferSizes {
     metal::uint size0;
 };
 
-struct type_3 {
+struct type_2 {
+    metal::float2x2 inner[2];
+};
+struct type_5 {
     metal::uint2 inner[2];
 };
-typedef int type_5[1];
+typedef int type_7[1];
 struct Bar {
     metal::float4x4 matrix;
+    type_2 matrix_array;
     metal::atomic_int atom;
-    char _pad2[4];
-    type_3 arr;
-    type_5 data;
+    char _pad3[4];
+    type_5 arr;
+    type_7 data;
 };
-struct type_11 {
+struct type_13 {
     int inner[5];
 };
 
@@ -39,19 +43,19 @@ vertex fooOutput foo(
 , constant _mslBufferSizes& _buffer_sizes [[buffer(24)]]
 ) {
     float foo_1 = 0.0;
-    type_11 c;
+    type_13 c;
     float baz = foo_1;
     foo_1 = 1.0;
     metal::float4x4 matrix = bar.matrix;
-    type_3 arr = bar.arr;
+    type_5 arr = bar.arr;
     float b = bar.matrix[3].x;
-    int a = bar.data[(1 + (_buffer_sizes.size0 - 88 - 4) / 8) - 2u];
+    int a = bar.data[(1 + (_buffer_sizes.size0 - 120 - 4) / 8) - 2u];
     float _e25 = read_from_private(foo_1);
     bar.matrix[1].z = 1.0;
     bar.matrix = metal::float4x4(metal::float4(0.0), metal::float4(1.0), metal::float4(2.0), metal::float4(3.0));
-    for(int _i=0; _i<2; ++_i) bar.arr.inner[_i] = type_3 {metal::uint2(0u), metal::uint2(1u)}.inner[_i];
+    for(int _i=0; _i<2; ++_i) bar.arr.inner[_i] = type_5 {metal::uint2(0u), metal::uint2(1u)}.inner[_i];
     bar.data[1] = 1;
-    for(int _i=0; _i<5; ++_i) c.inner[_i] = type_11 {a, static_cast<int>(b), 3, 4, 5}.inner[_i];
+    for(int _i=0; _i<5; ++_i) c.inner[_i] = type_13 {a, static_cast<int>(b), 3, 4, 5}.inner[_i];
     c.inner[vi + 1u] = 42;
     int value = c.inner[vi];
     return fooOutput { matrix * static_cast<metal::float4>(metal::int4(value)) };