From 0b87d1918a906da57341366de4d2b875e09753ee Mon Sep 17 00:00:00 2001 From: Erich Gubler Date: Mon, 20 Mar 2023 12:21:52 -0400 Subject: [PATCH] fix(hlsl): emit constructor functions for arrays (#2281) * test(hlsl-out): add failing case for array as ctor arg See issue #2184. * fix(hlsl): emit constructor functions for arrays --- src/back/hlsl/help.rs | 6 +++ src/back/hlsl/storage.rs | 8 +++- tests/in/array-in-ctor.param.ron | 2 + tests/in/array-in-ctor.wgsl | 10 +++++ .../glsl/array-in-ctor.cs_main.Compute.glsl | 17 +++++++++ tests/out/hlsl/access.hlsl | 14 +++---- tests/out/hlsl/array-in-ctor.hlsl | 24 ++++++++++++ tests/out/hlsl/array-in-ctor.hlsl.config | 3 ++ tests/out/msl/array-in-ctor.msl | 18 +++++++++ tests/out/spv/array-in-ctor.spvasm | 38 +++++++++++++++++++ tests/out/wgsl/array-in-ctor.wgsl | 11 ++++++ tests/snapshots.rs | 4 ++ 12 files changed, 146 insertions(+), 9 deletions(-) create mode 100644 tests/in/array-in-ctor.param.ron create mode 100644 tests/in/array-in-ctor.wgsl create mode 100644 tests/out/glsl/array-in-ctor.cs_main.Compute.glsl create mode 100644 tests/out/hlsl/array-in-ctor.hlsl create mode 100644 tests/out/hlsl/array-in-ctor.hlsl.config create mode 100644 tests/out/msl/array-in-ctor.msl create mode 100644 tests/out/spv/array-in-ctor.spvasm create mode 100644 tests/out/wgsl/array-in-ctor.wgsl diff --git a/src/back/hlsl/help.rs b/src/back/hlsl/help.rs index bb18bf3ab4..8ae9baf62d 100644 --- a/src/back/hlsl/help.rs +++ b/src/back/hlsl/help.rs @@ -882,6 +882,12 @@ impl<'a, W: Write> super::Writer<'a, W> { } crate::TypeInner::Array { base, .. } => { write_wrapped_constructor(writer, base, module)?; + let constructor = WrappedConstructor { ty }; + if !writer.wrapped.constructors.contains(&constructor) { + writer + .write_wrapped_constructor_function(module, constructor)?; + writer.wrapped.constructors.insert(constructor); + } } _ => {} }; diff --git a/src/back/hlsl/storage.rs b/src/back/hlsl/storage.rs index 21396c953a..647a6bf870 100644 --- a/src/back/hlsl/storage.rs +++ b/src/back/hlsl/storage.rs @@ -147,12 +147,16 @@ impl super::Writer<'_, W> { size: crate::ArraySize::Constant(const_handle), .. } => { - write!(self.out, "{{")?; + let constructor = super::help::WrappedConstructor { + ty: result_ty.handle().unwrap(), + }; + self.write_wrapped_constructor_function_name(module, constructor)?; + write!(self.out, "(")?; let count = module.constants[const_handle].to_array_length().unwrap(); let stride = module.types[base].inner.size(&module.constants); let iter = (0..count).map(|i| (TypeResolution::Handle(base), stride * i)); self.write_storage_load_sequence(module, var_handle, iter, func_ctx)?; - write!(self.out, "}}")?; + write!(self.out, ")")?; } crate::TypeInner::Struct { ref members, .. } => { let constructor = super::help::WrappedConstructor { diff --git a/tests/in/array-in-ctor.param.ron b/tests/in/array-in-ctor.param.ron new file mode 100644 index 0000000000..72873dd667 --- /dev/null +++ b/tests/in/array-in-ctor.param.ron @@ -0,0 +1,2 @@ +( +) diff --git a/tests/in/array-in-ctor.wgsl b/tests/in/array-in-ctor.wgsl new file mode 100644 index 0000000000..a7ca66be61 --- /dev/null +++ b/tests/in/array-in-ctor.wgsl @@ -0,0 +1,10 @@ +struct Ah { + inner: array, +}; +@group(0) @binding(0) +var ah: Ah; + +@compute @workgroup_size(1) +fn cs_main() { + _ = ah; +} diff --git a/tests/out/glsl/array-in-ctor.cs_main.Compute.glsl b/tests/out/glsl/array-in-ctor.cs_main.Compute.glsl new file mode 100644 index 0000000000..e8245d55ee --- /dev/null +++ b/tests/out/glsl/array-in-ctor.cs_main.Compute.glsl @@ -0,0 +1,17 @@ +#version 310 es + +precision highp float; +precision highp int; + +layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; + +struct Ah { + float inner[2]; +}; +layout(std430) readonly buffer Ah_block_0Compute { Ah _group_0_binding_0_cs; }; + + +void main() { + Ah unnamed = _group_0_binding_0_cs; +} + diff --git a/tests/out/hlsl/access.hlsl b/tests/out/hlsl/access.hlsl index a8d7d80c99..f20cb726db 100644 --- a/tests/out/hlsl/access.hlsl +++ b/tests/out/hlsl/access.hlsl @@ -249,6 +249,12 @@ void assign_array_through_ptr_fn(inout float4 foo_2[2]) return; } +typedef uint2 ret_Constructarray2_uint2_[2]; +ret_Constructarray2_uint2_ Constructarray2_uint2_(uint2 arg0, uint2 arg1) { + uint2 ret[2] = { arg0, arg1 }; + return ret; +} + uint NagaBufferLengthRW(RWByteAddressBuffer buffer) { uint ret; @@ -273,7 +279,7 @@ float4 foo_vert(uint vi : SV_VertexID) : SV_Position test_matrix_within_struct_accesses(); test_matrix_within_array_within_struct_accesses(); float4x3 _matrix = float4x3(asfloat(bar.Load3(0+0)), asfloat(bar.Load3(0+16)), asfloat(bar.Load3(0+32)), asfloat(bar.Load3(0+48))); - uint2 arr_1[2] = {asuint(bar.Load2(144+0)), asuint(bar.Load2(144+8))}; + uint2 arr_1[2] = Constructarray2_uint2_(asuint(bar.Load2(144+0)), asuint(bar.Load2(144+8))); float b = asfloat(bar.Load(0+48+0)); int a_1 = asint(bar.Load(0+(((NagaBufferLengthRW(bar) - 160) / 8) - 2u)*8+160)); int2 c = asint(qux.Load2(0)); @@ -285,12 +291,6 @@ float4 foo_vert(uint vi : SV_VertexID) : SV_Position return float4(mul(float4((value).xxxx), _matrix), 2.0); } -typedef uint2 ret_Constructarray2_uint2_[2]; -ret_Constructarray2_uint2_ Constructarray2_uint2_(uint2 arg0, uint2 arg1) { - uint2 ret[2] = { arg0, arg1 }; - return ret; -} - float4 foo_frag() : SV_Target0 { bar.Store(8+16+0, asuint(1.0)); diff --git a/tests/out/hlsl/array-in-ctor.hlsl b/tests/out/hlsl/array-in-ctor.hlsl new file mode 100644 index 0000000000..edaa954737 --- /dev/null +++ b/tests/out/hlsl/array-in-ctor.hlsl @@ -0,0 +1,24 @@ + +struct Ah { + float inner[2]; +}; + +ByteAddressBuffer ah : register(t0); + +typedef float ret_Constructarray2_float_[2]; +ret_Constructarray2_float_ Constructarray2_float_(float arg0, float arg1) { + float ret[2] = { arg0, arg1 }; + return ret; +} + +Ah ConstructAh(float arg0[2]) { + Ah ret = (Ah)0; + ret.inner = arg0; + return ret; +} + +[numthreads(1, 1, 1)] +void cs_main() +{ + Ah unnamed = ConstructAh(Constructarray2_float_(asfloat(ah.Load(0+0)), asfloat(ah.Load(0+4)))); +} diff --git a/tests/out/hlsl/array-in-ctor.hlsl.config b/tests/out/hlsl/array-in-ctor.hlsl.config new file mode 100644 index 0000000000..522d99ed1a --- /dev/null +++ b/tests/out/hlsl/array-in-ctor.hlsl.config @@ -0,0 +1,3 @@ +vertex=() +fragment=() +compute=(cs_main:cs_5_1 ) diff --git a/tests/out/msl/array-in-ctor.msl b/tests/out/msl/array-in-ctor.msl new file mode 100644 index 0000000000..eab190df75 --- /dev/null +++ b/tests/out/msl/array-in-ctor.msl @@ -0,0 +1,18 @@ +// language: metal2.0 +#include +#include + +using metal::uint; + +struct type_1 { + float inner[2]; +}; +struct Ah { + type_1 inner; +}; + +kernel void cs_main( + device Ah const& ah [[user(fake0)]] +) { + Ah unnamed = ah; +} diff --git a/tests/out/spv/array-in-ctor.spvasm b/tests/out/spv/array-in-ctor.spvasm new file mode 100644 index 0000000000..8ad438ce0a --- /dev/null +++ b/tests/out/spv/array-in-ctor.spvasm @@ -0,0 +1,38 @@ +; SPIR-V +; Version: 1.1 +; Generator: rspirv +; Bound: 20 +OpCapability Shader +OpExtension "SPV_KHR_storage_buffer_storage_class" +%1 = OpExtInstImport "GLSL.std.450" +OpMemoryModel Logical GLSL450 +OpEntryPoint GLCompute %12 "cs_main" +OpExecutionMode %12 LocalSize 1 1 1 +OpDecorate %6 ArrayStride 4 +OpMemberDecorate %7 0 Offset 0 +OpDecorate %8 NonWritable +OpDecorate %8 DescriptorSet 0 +OpDecorate %8 Binding 0 +OpDecorate %9 Block +OpMemberDecorate %9 0 Offset 0 +%2 = OpTypeVoid +%4 = OpTypeInt 32 1 +%3 = OpConstant %4 2 +%5 = OpTypeFloat 32 +%6 = OpTypeArray %5 %3 +%7 = OpTypeStruct %6 +%9 = OpTypeStruct %7 +%10 = OpTypePointer StorageBuffer %9 +%8 = OpVariable %10 StorageBuffer +%13 = OpTypeFunction %2 +%14 = OpTypePointer StorageBuffer %7 +%16 = OpTypeInt 32 0 +%15 = OpConstant %16 0 +%12 = OpFunction %2 None %13 +%11 = OpLabel +%17 = OpAccessChain %14 %8 %15 +OpBranch %18 +%18 = OpLabel +%19 = OpLoad %7 %17 +OpReturn +OpFunctionEnd \ No newline at end of file diff --git a/tests/out/wgsl/array-in-ctor.wgsl b/tests/out/wgsl/array-in-ctor.wgsl new file mode 100644 index 0000000000..f4a8f4ddd1 --- /dev/null +++ b/tests/out/wgsl/array-in-ctor.wgsl @@ -0,0 +1,11 @@ +struct Ah { + inner: array, +} + +@group(0) @binding(0) +var ah: Ah; + +@compute @workgroup_size(1, 1, 1) +fn cs_main() { + _ = ah; +} diff --git a/tests/snapshots.rs b/tests/snapshots.rs index 16be39589f..691c074a93 100644 --- a/tests/snapshots.rs +++ b/tests/snapshots.rs @@ -426,6 +426,10 @@ fn convert_wgsl() { let root = env!("CARGO_MANIFEST_DIR"); let inputs = [ + ( + "array-in-ctor", + Targets::SPIRV | Targets::METAL | Targets::GLSL | Targets::HLSL | Targets::WGSL, + ), ( "empty", Targets::SPIRV | Targets::METAL | Targets::GLSL | Targets::HLSL | Targets::WGSL,