From 9f8338cda784a77b68247c384b5a9515472d51e0 Mon Sep 17 00:00:00 2001 From: Dzmitry Malyshau Date: Sun, 26 Dec 2021 19:23:20 -0500 Subject: [PATCH] hlsl: support arrays of matrices --- src/back/hlsl/writer.rs | 19 +- tests/in/access.wgsl | 1 + tests/out/glsl/access.atomics.Compute.glsl | 1 + tests/out/glsl/access.foo.Vertex.glsl | 1 + tests/out/hlsl/access.hlsl | 30 +-- tests/out/msl/access.msl | 26 +- tests/out/spv/access.spvasm | 297 +++++++++++---------- tests/out/wgsl/access.wgsl | 1 + 8 files changed, 191 insertions(+), 185 deletions(-) diff --git a/src/back/hlsl/writer.rs b/src/back/hlsl/writer.rs index 0280b50585..c9e8880e3a 100644 --- a/src/back/hlsl/writer.rs +++ b/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, " {}", diff --git a/tests/in/access.wgsl b/tests/in/access.wgsl index e8ef082fbc..ed088e7de1 100644 --- a/tests/in/access.wgsl +++ b/tests/in/access.wgsl @@ -2,6 +2,7 @@ struct Bar { matrix: mat4x4; + matrix_array: array, 2>; atom: atomic; arr: [[stride(8)]] array, 2>; data: [[stride(8)]] array; diff --git a/tests/out/glsl/access.atomics.Compute.glsl b/tests/out/glsl/access.atomics.Compute.glsl index fbc36cfced..cca755cd68 100644 --- a/tests/out/glsl/access.atomics.Compute.glsl +++ b/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[]; diff --git a/tests/out/glsl/access.foo.Vertex.glsl b/tests/out/glsl/access.foo.Vertex.glsl index 2f080d9f0c..9b24f67403 100644 --- a/tests/out/glsl/access.foo.Vertex.glsl +++ b/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[]; diff --git a/tests/out/hlsl/access.hlsl b/tests/out/hlsl/access.hlsl index 91b15a93c5..aa89e7f901 100644 --- a/tests/out/hlsl/access.hlsl +++ b/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; } diff --git a/tests/out/msl/access.msl b/tests/out/msl/access.msl index fda0a32019..53897e1adf 100644 --- a/tests/out/msl/access.msl +++ b/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(b), 3, 4, 5}.inner[_i]; + for(int _i=0; _i<5; ++_i) c.inner[_i] = type_13 {a, static_cast(b), 3, 4, 5}.inner[_i]; c.inner[vi + 1u] = 42; int value = c.inner[vi]; return fooOutput { matrix * static_cast(metal::int4(value)) }; diff --git a/tests/out/spv/access.spvasm b/tests/out/spv/access.spvasm index cd6fea5375..6a599ad51a 100644 --- a/tests/out/spv/access.spvasm +++ b/tests/out/spv/access.spvasm @@ -1,43 +1,48 @@ ; SPIR-V ; Version: 1.1 ; Generator: rspirv -; Bound: 115 +; Bound: 119 OpCapability Shader OpExtension "SPV_KHR_storage_buffer_storage_class" %1 = OpExtInstImport "GLSL.std.450" OpMemoryModel Logical GLSL450 -OpEntryPoint Vertex %47 "foo" %42 %45 -OpEntryPoint GLCompute %92 "atomics" -OpExecutionMode %92 LocalSize 1 1 1 +OpEntryPoint Vertex %50 "foo" %45 %48 +OpEntryPoint GLCompute %96 "atomics" +OpExecutionMode %96 LocalSize 1 1 1 OpSource GLSL 450 -OpMemberName %26 0 "matrix" -OpMemberName %26 1 "atom" -OpMemberName %26 2 "arr" -OpMemberName %26 3 "data" -OpName %26 "Bar" -OpName %30 "bar" -OpName %33 "foo" -OpName %34 "read_from_private" -OpName %38 "foo" -OpName %39 "c" -OpName %42 "vi" -OpName %47 "foo" -OpName %90 "tmp" -OpName %92 "atomics" -OpDecorate %24 ArrayStride 8 -OpDecorate %25 ArrayStride 8 -OpMemberDecorate %26 0 Offset 0 -OpMemberDecorate %26 0 ColMajor -OpMemberDecorate %26 0 MatrixStride 16 -OpMemberDecorate %26 1 Offset 64 -OpMemberDecorate %26 2 Offset 72 -OpMemberDecorate %26 3 Offset 88 -OpDecorate %29 ArrayStride 4 -OpDecorate %30 DescriptorSet 0 -OpDecorate %30 Binding 0 -OpDecorate %26 Block -OpDecorate %42 BuiltIn VertexIndex -OpDecorate %45 BuiltIn Position +OpMemberName %29 0 "matrix" +OpMemberName %29 1 "matrix_array" +OpMemberName %29 2 "atom" +OpMemberName %29 3 "arr" +OpMemberName %29 4 "data" +OpName %29 "Bar" +OpName %33 "bar" +OpName %36 "foo" +OpName %37 "read_from_private" +OpName %41 "foo" +OpName %42 "c" +OpName %45 "vi" +OpName %50 "foo" +OpName %94 "tmp" +OpName %96 "atomics" +OpDecorate %25 ArrayStride 16 +OpDecorate %27 ArrayStride 8 +OpDecorate %28 ArrayStride 8 +OpMemberDecorate %29 0 Offset 0 +OpMemberDecorate %29 0 ColMajor +OpMemberDecorate %29 0 MatrixStride 16 +OpMemberDecorate %29 1 Offset 64 +OpMemberDecorate %29 1 ColMajor +OpMemberDecorate %29 1 MatrixStride 8 +OpMemberDecorate %29 2 Offset 96 +OpMemberDecorate %29 3 Offset 104 +OpMemberDecorate %29 4 Offset 120 +OpDecorate %32 ArrayStride 4 +OpDecorate %33 DescriptorSet 0 +OpDecorate %33 Binding 0 +OpDecorate %29 Block +OpDecorate %45 BuiltIn VertexIndex +OpDecorate %48 BuiltIn Position %2 = OpTypeVoid %4 = OpTypeInt 32 1 %3 = OpConstant %4 2 @@ -59,121 +64,125 @@ OpDecorate %45 BuiltIn Position %20 = OpConstant %4 42 %22 = OpTypeVector %6 4 %21 = OpTypeMatrix %22 4 -%23 = OpTypeVector %9 2 -%24 = OpTypeArray %23 %3 -%25 = OpTypeRuntimeArray %4 -%26 = OpTypeStruct %21 %4 %24 %25 -%27 = OpTypePointer Function %6 -%28 = OpTypePointer StorageBuffer %4 -%29 = OpTypeArray %4 %17 -%31 = OpTypePointer StorageBuffer %26 -%30 = OpVariable %31 StorageBuffer -%35 = OpTypeFunction %6 %27 -%40 = OpTypePointer Function %29 -%43 = OpTypePointer Input %9 -%42 = OpVariable %43 Input -%46 = OpTypePointer Output %22 -%45 = OpVariable %46 Output -%48 = OpTypeFunction %2 -%51 = OpTypePointer StorageBuffer %21 -%54 = OpTypePointer StorageBuffer %24 -%57 = OpTypePointer StorageBuffer %22 -%58 = OpTypePointer StorageBuffer %6 -%61 = OpTypePointer StorageBuffer %25 -%82 = OpTypePointer Function %4 -%86 = OpTypeVector %4 4 -%94 = OpTypePointer StorageBuffer %4 -%97 = OpConstant %9 64 -%34 = OpFunction %6 None %35 -%33 = OpFunctionParameter %27 -%32 = OpLabel -OpBranch %36 -%36 = OpLabel -%37 = OpLoad %6 %33 -OpReturnValue %37 +%24 = OpTypeVector %6 2 +%23 = OpTypeMatrix %24 2 +%25 = OpTypeArray %23 %3 +%26 = OpTypeVector %9 2 +%27 = OpTypeArray %26 %3 +%28 = OpTypeRuntimeArray %4 +%29 = OpTypeStruct %21 %25 %4 %27 %28 +%30 = OpTypePointer Function %6 +%31 = OpTypePointer StorageBuffer %4 +%32 = OpTypeArray %4 %17 +%34 = OpTypePointer StorageBuffer %29 +%33 = OpVariable %34 StorageBuffer +%38 = OpTypeFunction %6 %30 +%43 = OpTypePointer Function %32 +%46 = OpTypePointer Input %9 +%45 = OpVariable %46 Input +%49 = OpTypePointer Output %22 +%48 = OpVariable %49 Output +%51 = OpTypeFunction %2 +%54 = OpTypePointer StorageBuffer %21 +%57 = OpTypePointer StorageBuffer %27 +%60 = OpTypePointer StorageBuffer %22 +%61 = OpTypePointer StorageBuffer %6 +%64 = OpTypePointer StorageBuffer %28 +%67 = OpConstant %9 4 +%86 = OpTypePointer Function %4 +%90 = OpTypeVector %4 4 +%98 = OpTypePointer StorageBuffer %4 +%101 = OpConstant %9 64 +%37 = OpFunction %6 None %38 +%36 = OpFunctionParameter %30 +%35 = OpLabel +OpBranch %39 +%39 = OpLabel +%40 = OpLoad %6 %36 +OpReturnValue %40 OpFunctionEnd -%47 = OpFunction %2 None %48 -%41 = OpLabel -%38 = OpVariable %27 Function %5 -%39 = OpVariable %40 Function -%44 = OpLoad %9 %42 -OpBranch %49 -%49 = OpLabel -%50 = OpLoad %6 %38 -OpStore %38 %7 -%52 = OpAccessChain %51 %30 %15 -%53 = OpLoad %21 %52 -%55 = OpAccessChain %54 %30 %10 -%56 = OpLoad %24 %55 -%59 = OpAccessChain %58 %30 %15 %8 %15 -%60 = OpLoad %6 %59 -%62 = OpArrayLength %9 %30 3 -%63 = OpISub %9 %62 %10 -%64 = OpAccessChain %28 %30 %8 %63 -%65 = OpLoad %4 %64 -%66 = OpFunctionCall %6 %34 %38 -%67 = OpAccessChain %58 %30 %15 %16 %10 -OpStore %67 %7 -%68 = OpCompositeConstruct %22 %5 %5 %5 %5 -%69 = OpCompositeConstruct %22 %7 %7 %7 %7 -%70 = OpCompositeConstruct %22 %13 %13 %13 %13 -%71 = OpCompositeConstruct %22 %14 %14 %14 %14 -%72 = OpCompositeConstruct %21 %68 %69 %70 %71 -%73 = OpAccessChain %51 %30 %15 -OpStore %73 %72 -%74 = OpCompositeConstruct %23 %15 %15 -%75 = OpCompositeConstruct %23 %16 %16 -%76 = OpCompositeConstruct %24 %74 %75 -%77 = OpAccessChain %54 %30 %10 +%50 = OpFunction %2 None %51 +%44 = OpLabel +%41 = OpVariable %30 Function %5 +%42 = OpVariable %43 Function +%47 = OpLoad %9 %45 +OpBranch %52 +%52 = OpLabel +%53 = OpLoad %6 %41 +OpStore %41 %7 +%55 = OpAccessChain %54 %33 %15 +%56 = OpLoad %21 %55 +%58 = OpAccessChain %57 %33 %8 +%59 = OpLoad %27 %58 +%62 = OpAccessChain %61 %33 %15 %8 %15 +%63 = OpLoad %6 %62 +%65 = OpArrayLength %9 %33 4 +%66 = OpISub %9 %65 %10 +%68 = OpAccessChain %31 %33 %67 %66 +%69 = OpLoad %4 %68 +%70 = OpFunctionCall %6 %37 %41 +%71 = OpAccessChain %61 %33 %15 %16 %10 +OpStore %71 %7 +%72 = OpCompositeConstruct %22 %5 %5 %5 %5 +%73 = OpCompositeConstruct %22 %7 %7 %7 %7 +%74 = OpCompositeConstruct %22 %13 %13 %13 %13 +%75 = OpCompositeConstruct %22 %14 %14 %14 %14 +%76 = OpCompositeConstruct %21 %72 %73 %74 %75 +%77 = OpAccessChain %54 %33 %15 OpStore %77 %76 -%78 = OpAccessChain %28 %30 %8 %16 -OpStore %78 %12 -%79 = OpConvertFToS %4 %60 -%80 = OpCompositeConstruct %29 %65 %79 %18 %19 %17 -OpStore %39 %80 -%81 = OpIAdd %9 %44 %16 -%83 = OpAccessChain %82 %39 %81 -OpStore %83 %20 -%84 = OpAccessChain %82 %39 %44 -%85 = OpLoad %4 %84 -%87 = OpCompositeConstruct %86 %85 %85 %85 %85 -%88 = OpConvertSToF %22 %87 -%89 = OpMatrixTimesVector %22 %53 %88 -OpStore %45 %89 +%78 = OpCompositeConstruct %26 %15 %15 +%79 = OpCompositeConstruct %26 %16 %16 +%80 = OpCompositeConstruct %27 %78 %79 +%81 = OpAccessChain %57 %33 %8 +OpStore %81 %80 +%82 = OpAccessChain %31 %33 %67 %16 +OpStore %82 %12 +%83 = OpConvertFToS %4 %63 +%84 = OpCompositeConstruct %32 %69 %83 %18 %19 %17 +OpStore %42 %84 +%85 = OpIAdd %9 %47 %16 +%87 = OpAccessChain %86 %42 %85 +OpStore %87 %20 +%88 = OpAccessChain %86 %42 %47 +%89 = OpLoad %4 %88 +%91 = OpCompositeConstruct %90 %89 %89 %89 %89 +%92 = OpConvertSToF %22 %91 +%93 = OpMatrixTimesVector %22 %56 %92 +OpStore %48 %93 OpReturn OpFunctionEnd -%92 = OpFunction %2 None %48 -%91 = OpLabel -%90 = OpVariable %82 Function -OpBranch %93 -%93 = OpLabel -%95 = OpAccessChain %94 %30 %16 -%96 = OpAtomicLoad %4 %95 %12 %97 -%99 = OpAccessChain %94 %30 %16 -%98 = OpAtomicIAdd %4 %99 %12 %97 %17 -OpStore %90 %98 -%101 = OpAccessChain %94 %30 %16 -%100 = OpAtomicISub %4 %101 %12 %97 %17 -OpStore %90 %100 -%103 = OpAccessChain %94 %30 %16 -%102 = OpAtomicAnd %4 %103 %12 %97 %17 -OpStore %90 %102 -%105 = OpAccessChain %94 %30 %16 -%104 = OpAtomicOr %4 %105 %12 %97 %17 -OpStore %90 %104 -%107 = OpAccessChain %94 %30 %16 -%106 = OpAtomicXor %4 %107 %12 %97 %17 -OpStore %90 %106 -%109 = OpAccessChain %94 %30 %16 -%108 = OpAtomicSMin %4 %109 %12 %97 %17 -OpStore %90 %108 -%111 = OpAccessChain %94 %30 %16 -%110 = OpAtomicSMax %4 %111 %12 %97 %17 -OpStore %90 %110 -%113 = OpAccessChain %94 %30 %16 -%112 = OpAtomicExchange %4 %113 %12 %97 %17 -OpStore %90 %112 -%114 = OpAccessChain %94 %30 %16 -OpAtomicStore %114 %12 %97 %96 +%96 = OpFunction %2 None %51 +%95 = OpLabel +%94 = OpVariable %86 Function +OpBranch %97 +%97 = OpLabel +%99 = OpAccessChain %98 %33 %10 +%100 = OpAtomicLoad %4 %99 %12 %101 +%103 = OpAccessChain %98 %33 %10 +%102 = OpAtomicIAdd %4 %103 %12 %101 %17 +OpStore %94 %102 +%105 = OpAccessChain %98 %33 %10 +%104 = OpAtomicISub %4 %105 %12 %101 %17 +OpStore %94 %104 +%107 = OpAccessChain %98 %33 %10 +%106 = OpAtomicAnd %4 %107 %12 %101 %17 +OpStore %94 %106 +%109 = OpAccessChain %98 %33 %10 +%108 = OpAtomicOr %4 %109 %12 %101 %17 +OpStore %94 %108 +%111 = OpAccessChain %98 %33 %10 +%110 = OpAtomicXor %4 %111 %12 %101 %17 +OpStore %94 %110 +%113 = OpAccessChain %98 %33 %10 +%112 = OpAtomicSMin %4 %113 %12 %101 %17 +OpStore %94 %112 +%115 = OpAccessChain %98 %33 %10 +%114 = OpAtomicSMax %4 %115 %12 %101 %17 +OpStore %94 %114 +%117 = OpAccessChain %98 %33 %10 +%116 = OpAtomicExchange %4 %117 %12 %101 %17 +OpStore %94 %116 +%118 = OpAccessChain %98 %33 %10 +OpAtomicStore %118 %12 %101 %100 OpReturn OpFunctionEnd \ No newline at end of file diff --git a/tests/out/wgsl/access.wgsl b/tests/out/wgsl/access.wgsl index e437077157..ca5313a808 100644 --- a/tests/out/wgsl/access.wgsl +++ b/tests/out/wgsl/access.wgsl @@ -1,5 +1,6 @@ struct Bar { matrix: mat4x4; + matrix_array: [[stride(16)]] array,2>; atom: atomic; arr: [[stride(8)]] array,2>; data: [[stride(8)]] array;