From f3e810b8b30cc9a4bab772fac6ea1f7031b37464 Mon Sep 17 00:00:00 2001 From: Hans-Kristian Arntzen Date: Wed, 2 May 2018 09:38:41 +0200 Subject: [PATCH] Fix OpCompositeConstruct with arrays in MSL. --- .../shaders-msl/comp/composite-construct.comp | 15 ++++ .../shaders-msl/comp/composite-construct.comp | 69 +++++++++++++++++++ shaders-msl/comp/composite-construct.comp | 40 +++++++++++ spirv_glsl.cpp | 2 +- 4 files changed, 125 insertions(+), 1 deletion(-) create mode 100644 reference/opt/shaders-msl/comp/composite-construct.comp create mode 100644 reference/shaders-msl/comp/composite-construct.comp create mode 100644 shaders-msl/comp/composite-construct.comp diff --git a/reference/opt/shaders-msl/comp/composite-construct.comp b/reference/opt/shaders-msl/comp/composite-construct.comp new file mode 100644 index 0000000..bbf4b13 --- /dev/null +++ b/reference/opt/shaders-msl/comp/composite-construct.comp @@ -0,0 +1,15 @@ +#include +#include + +using namespace metal; + +struct SSBO0 +{ + float4 as[1]; +}; + +kernel void main0(device SSBO0& _41 [[buffer(0)]], device SSBO0& _55 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) +{ + _41.as[gl_GlobalInvocationID.x] = ((_41.as[gl_GlobalInvocationID.x] + _55.as[gl_GlobalInvocationID.x]) + _55.as[gl_GlobalInvocationID.x]) + float4(10.0); +} + diff --git a/reference/shaders-msl/comp/composite-construct.comp b/reference/shaders-msl/comp/composite-construct.comp new file mode 100644 index 0000000..5e748ef --- /dev/null +++ b/reference/shaders-msl/comp/composite-construct.comp @@ -0,0 +1,69 @@ +#pragma clang diagnostic ignored "-Wmissing-prototypes" + +#include +#include + +using namespace metal; + +struct SSBO0 +{ + float4 as[1]; +}; + +struct SSBO1 +{ + float4 bs[1]; +}; + +struct Composite +{ + float4 a[2]; + float4 b[2]; +}; + +constant float4 _66[2] = {float4(10.0), float4(30.0)}; +constant float _91[3] = {1.0, 1.0, 1.0}; +constant float _93[3] = {2.0, 2.0, 2.0}; +constant float _94[2][3] = {{1.0, 1.0, 1.0}, {2.0, 2.0, 2.0}}; + +// Implementation of an array copy function to cover GLSL's ability to copy an array via assignment. +template +void spvArrayCopy(thread T (&dst)[N], thread const T (&src)[N]) +{ + for (uint i = 0; i < N; dst[i] = src[i], i++); +} + +// An overload for constant arrays. +template +void spvArrayCopyConstant(thread T (&dst)[N], constant T (&src)[N]) +{ + for (uint i = 0; i < N; dst[i] = src[i], i++); +} + +float4 summe(thread const float4 (&values)[3][2]) +{ + return ((values[0][0] + values[2][1]) + values[0][1]) + values[1][0]; +} + +kernel void main0(device SSBO0& _41 [[buffer(0)]], device SSBO1& _55 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) +{ + float4 _60[2] = { _41.as[gl_GlobalInvocationID.x], _55.bs[gl_GlobalInvocationID.x] }; + float4 values[2]; + spvArrayCopy(values, _60); + float4 const_values[2] = {float4(10.0), float4(30.0)}; + float4 copy_values[2]; + spvArrayCopy(copy_values, const_values); + float4 copy_values2[2]; + spvArrayCopy(copy_values2, values); + float4 _76[3][2] = { values, copy_values, copy_values2 }; + float4 param[3][2]; + spvArrayCopy(param, _76); + _41.as[gl_GlobalInvocationID.x] = summe(param); + Composite c = Composite{ values, copy_values }; + float arrayofarray[2][3] = {{1.0, 1.0, 1.0}, {2.0, 2.0, 2.0}}; + float b = 10.0; + float _105[4] = { b, b, b, b }; + float values_scalar[4]; + spvArrayCopy(values_scalar, _105); +} + diff --git a/shaders-msl/comp/composite-construct.comp b/shaders-msl/comp/composite-construct.comp new file mode 100644 index 0000000..859c56f --- /dev/null +++ b/shaders-msl/comp/composite-construct.comp @@ -0,0 +1,40 @@ +#version 310 es +layout(local_size_x = 1) in; + +layout(std430, binding = 0) buffer SSBO0 +{ + vec4 as[]; +}; + +layout(std430, binding = 1) buffer SSBO1 +{ + vec4 bs[]; +}; + +vec4 summe(vec4 values[3][2]) +{ + return values[0][0] + values[2][1] + values[0][1] + values[1][0]; +} + +struct Composite +{ + vec4 a[2]; + vec4 b[2]; +}; + +void main() +{ + vec4 values[2] = vec4[](as[gl_GlobalInvocationID.x], bs[gl_GlobalInvocationID.x]); + vec4 const_values[2] = vec4[](vec4(10.0), vec4(30.0)); + vec4 copy_values[2]; + copy_values = const_values; + vec4 copy_values2[2] = values; + as[gl_GlobalInvocationID.x] = summe(vec4[][](values, copy_values, copy_values2)); + + Composite c = Composite(values, copy_values); + + float arrayofarray[2][3] = float[][](float[](1.0, 1.0, 1.0), float[](2.0, 2.0, 2.0)); + + float b = 10.0; + float values_scalar[4] = float[](b, b, b, b); +} diff --git a/spirv_glsl.cpp b/spirv_glsl.cpp index d064215..e5e1c52 100644 --- a/spirv_glsl.cpp +++ b/spirv_glsl.cpp @@ -6289,7 +6289,7 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction) { // Only use this path if we are building composites. // This path cannot be used for arithmetic. - if (backend.use_typed_initializer_list) + if (backend.use_typed_initializer_list && out_type.basetype == SPIRType::Struct) constructor_op += type_to_glsl_constructor(get(result_type)); constructor_op += "{ "; if (type_is_empty(out_type) && !backend.supports_empty_struct)