From 8cf99e7d44fd452018c76bf7596b1b8f58d511a5 Mon Sep 17 00:00:00 2001 From: Chip Davis Date: Wed, 26 Oct 2022 13:31:10 -0700 Subject: [PATCH] MSL: Implement `CompositeInsert` `OpSpecConstantOp`. This op creates a new composite constant with one element replaced. So, we reconstruct the `SPIRConstant` for the composite constant, but with one of the IDs replaced. Constant initializer lists are memoized for when the result of a `CompositeInsert` is used in another `CompositeInsert`. (I wanted to add a test case for GLSL as well, but for two things: 1. `glslang` in Vulkan mode chokes on the first constant array, insisting that its initializer needs to be a constant. [Bug in glslang?] 2. The declarations for the buffers used by the shader aren't emitted, regardless of whether Vulkan mode is enabled.) Fixes five tests under `dEQP-VK.spirv_assembly.instruction.*.opspecconstantop.vector_related`. --- ...p-spec-constant-op-vector-related.asm.comp | 111 ++++++++++++++ ...p-spec-constant-op-vector-related.asm.comp | 112 ++++++++++++++ ...p-spec-constant-op-vector-related.asm.comp | 107 +++++++++++++ spirv_cross_containers.hpp | 11 +- spirv_glsl.cpp | 142 ++++++++++++++++-- spirv_glsl.hpp | 6 + 6 files changed, 475 insertions(+), 14 deletions(-) create mode 100644 reference/opt/shaders-msl/asm/comp/op-spec-constant-op-vector-related.asm.comp create mode 100644 reference/shaders-msl/asm/comp/op-spec-constant-op-vector-related.asm.comp create mode 100644 shaders-msl/asm/comp/op-spec-constant-op-vector-related.asm.comp diff --git a/reference/opt/shaders-msl/asm/comp/op-spec-constant-op-vector-related.asm.comp b/reference/opt/shaders-msl/asm/comp/op-spec-constant-op-vector-related.asm.comp new file mode 100644 index 00000000..7c3eba06 --- /dev/null +++ b/reference/opt/shaders-msl/asm/comp/op-spec-constant-op-vector-related.asm.comp @@ -0,0 +1,111 @@ +#pragma clang diagnostic ignored "-Wmissing-prototypes" +#pragma clang diagnostic ignored "-Wmissing-braces" + +#include +#include + +using namespace metal; + +template +struct spvUnsafeArray +{ + T elements[Num ? Num : 1]; + + thread T& operator [] (size_t pos) thread + { + return elements[pos]; + } + constexpr const thread T& operator [] (size_t pos) const thread + { + return elements[pos]; + } + + device T& operator [] (size_t pos) device + { + return elements[pos]; + } + constexpr const device T& operator [] (size_t pos) const device + { + return elements[pos]; + } + + constexpr const constant T& operator [] (size_t pos) const constant + { + return elements[pos]; + } + + threadgroup T& operator [] (size_t pos) threadgroup + { + return elements[pos]; + } + constexpr const threadgroup T& operator [] (size_t pos) const threadgroup + { + return elements[pos]; + } +}; + +struct _29 +{ + spvUnsafeArray, 3> _m0; +}; + +struct _7 +{ + int _m0[1]; +}; + +constant int _3_tmp [[function_constant(0)]]; +constant int _3 = is_function_constant_defined(_3_tmp) ? _3_tmp : 0; +constant int _4_tmp [[function_constant(1)]]; +constant int _4 = is_function_constant_defined(_4_tmp) ? _4_tmp : 0; +constant int _5_tmp [[function_constant(2)]]; +constant int _5 = is_function_constant_defined(_5_tmp) ? _5_tmp : 0; +constant spvUnsafeArray _36 = spvUnsafeArray({ _3, 0, 0 }); +constant spvUnsafeArray _37 = spvUnsafeArray({ _3, _4, 0 }); +constant spvUnsafeArray _38 = spvUnsafeArray({ _3, _4, _5 }); +constant spvUnsafeArray _39 = spvUnsafeArray({ _4, 0, 0 }); +constant spvUnsafeArray _40 = spvUnsafeArray({ _4, _5, 0 }); +constant spvUnsafeArray _41 = spvUnsafeArray({ _4, _5, _3 }); +constant spvUnsafeArray _42 = spvUnsafeArray({ _5, 0, 0 }); +constant spvUnsafeArray _43 = spvUnsafeArray({ _5, _3, 0 }); +constant spvUnsafeArray _44 = spvUnsafeArray({ _5, _3, _4 }); +constant spvUnsafeArray, 3> _45 = spvUnsafeArray, 3>({ spvUnsafeArray({ _3, _4, _5 }), spvUnsafeArray({ 0, 0, 0 }), spvUnsafeArray({ 0, 0, 0 }) }); +constant spvUnsafeArray, 3> _46 = spvUnsafeArray, 3>({ spvUnsafeArray({ _3, _4, _5 }), spvUnsafeArray({ _4, _5, _3 }), spvUnsafeArray({ 0, 0, 0 }) }); +constant spvUnsafeArray, 3> _47 = spvUnsafeArray, 3>({ spvUnsafeArray({ _3, _4, _5 }), spvUnsafeArray({ _4, _5, _3 }), spvUnsafeArray({ _5, _3, _4 }) }); +constant _29 _48 = _29{ spvUnsafeArray, 3>({ spvUnsafeArray({ _3, _4, _5 }), spvUnsafeArray({ _4, _5, _3 }), spvUnsafeArray({ _5, _3, _4 }) }) }; +constant int _50 = _48._m0[0][0]; +constant int _51 = _48._m0[1][0]; +constant int _52 = _48._m0[0][1]; +constant int _53 = _48._m0[2][2]; +constant int _54 = _48._m0[2][0]; +constant int _55 = _48._m0[1][1]; +constant bool _56 = (_50 == _51); +constant bool _57 = (_52 == _53); +constant bool _58 = (_54 == _55); +constant int _59 = int(_56); +constant int _60 = int(_57); +constant int _61 = _58 ? 2 : 1; +constant int3 _62 = int3(_3, 0, 0); +constant int3 _63 = int3(0, _4, 0); +constant int3 _64 = int3(0, 0, _5); +constant int3 _65 = int3(_62.x, 0, _62.z); +constant int3 _66 = int3(0, _63.y, _63.x); +constant int3 _67 = int3(_64.z, 0, _64.z); +constant int3 _68 = int3(_65.y, _65.x, _66.y); +constant int3 _69 = int3(_67.z, _68.y, _68.z); +constant int _70 = _69.x; +constant int _71 = _69.y; +constant int _72 = _69.z; +constant int _73 = (_70 - _71); +constant int _74 = (_73 * _72); + +constant spvUnsafeArray _33 = spvUnsafeArray({ 0, 0, 0 }); +constant spvUnsafeArray, 3> _34 = spvUnsafeArray, 3>({ spvUnsafeArray({ 0, 0, 0 }), spvUnsafeArray({ 0, 0, 0 }), spvUnsafeArray({ 0, 0, 0 }) }); + +constant int3 _32 = {}; + +kernel void main0(device _7& _8 [[buffer(0)]], device _7& _9 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) +{ + _9._m0[gl_GlobalInvocationID.x] = _8._m0[gl_GlobalInvocationID.x] + ((((1 - _59) * _60) * (_61 - 1)) * _74); +} + diff --git a/reference/shaders-msl/asm/comp/op-spec-constant-op-vector-related.asm.comp b/reference/shaders-msl/asm/comp/op-spec-constant-op-vector-related.asm.comp new file mode 100644 index 00000000..050dca4a --- /dev/null +++ b/reference/shaders-msl/asm/comp/op-spec-constant-op-vector-related.asm.comp @@ -0,0 +1,112 @@ +#pragma clang diagnostic ignored "-Wmissing-prototypes" +#pragma clang diagnostic ignored "-Wmissing-braces" + +#include +#include + +using namespace metal; + +template +struct spvUnsafeArray +{ + T elements[Num ? Num : 1]; + + thread T& operator [] (size_t pos) thread + { + return elements[pos]; + } + constexpr const thread T& operator [] (size_t pos) const thread + { + return elements[pos]; + } + + device T& operator [] (size_t pos) device + { + return elements[pos]; + } + constexpr const device T& operator [] (size_t pos) const device + { + return elements[pos]; + } + + constexpr const constant T& operator [] (size_t pos) const constant + { + return elements[pos]; + } + + threadgroup T& operator [] (size_t pos) threadgroup + { + return elements[pos]; + } + constexpr const threadgroup T& operator [] (size_t pos) const threadgroup + { + return elements[pos]; + } +}; + +struct _29 +{ + spvUnsafeArray, 3> _m0; +}; + +struct _7 +{ + int _m0[1]; +}; + +constant int _3_tmp [[function_constant(0)]]; +constant int _3 = is_function_constant_defined(_3_tmp) ? _3_tmp : 0; +constant int _4_tmp [[function_constant(1)]]; +constant int _4 = is_function_constant_defined(_4_tmp) ? _4_tmp : 0; +constant int _5_tmp [[function_constant(2)]]; +constant int _5 = is_function_constant_defined(_5_tmp) ? _5_tmp : 0; +constant spvUnsafeArray _36 = spvUnsafeArray({ _3, 0, 0 }); +constant spvUnsafeArray _37 = spvUnsafeArray({ _3, _4, 0 }); +constant spvUnsafeArray _38 = spvUnsafeArray({ _3, _4, _5 }); +constant spvUnsafeArray _39 = spvUnsafeArray({ _4, 0, 0 }); +constant spvUnsafeArray _40 = spvUnsafeArray({ _4, _5, 0 }); +constant spvUnsafeArray _41 = spvUnsafeArray({ _4, _5, _3 }); +constant spvUnsafeArray _42 = spvUnsafeArray({ _5, 0, 0 }); +constant spvUnsafeArray _43 = spvUnsafeArray({ _5, _3, 0 }); +constant spvUnsafeArray _44 = spvUnsafeArray({ _5, _3, _4 }); +constant spvUnsafeArray, 3> _45 = spvUnsafeArray, 3>({ spvUnsafeArray({ _3, _4, _5 }), spvUnsafeArray({ 0, 0, 0 }), spvUnsafeArray({ 0, 0, 0 }) }); +constant spvUnsafeArray, 3> _46 = spvUnsafeArray, 3>({ spvUnsafeArray({ _3, _4, _5 }), spvUnsafeArray({ _4, _5, _3 }), spvUnsafeArray({ 0, 0, 0 }) }); +constant spvUnsafeArray, 3> _47 = spvUnsafeArray, 3>({ spvUnsafeArray({ _3, _4, _5 }), spvUnsafeArray({ _4, _5, _3 }), spvUnsafeArray({ _5, _3, _4 }) }); +constant _29 _48 = _29{ spvUnsafeArray, 3>({ spvUnsafeArray({ _3, _4, _5 }), spvUnsafeArray({ _4, _5, _3 }), spvUnsafeArray({ _5, _3, _4 }) }) }; +constant _29 _49 = _29{ spvUnsafeArray, 3>({ spvUnsafeArray({ _3, _4, _5 }), spvUnsafeArray({ _4, _5, _5 }), spvUnsafeArray({ _5, _3, _4 }) }) }; +constant int _50 = _48._m0[0][0]; +constant int _51 = _48._m0[1][0]; +constant int _52 = _48._m0[0][1]; +constant int _53 = _48._m0[2][2]; +constant int _54 = _48._m0[2][0]; +constant int _55 = _48._m0[1][1]; +constant bool _56 = (_50 == _51); +constant bool _57 = (_52 == _53); +constant bool _58 = (_54 == _55); +constant int _59 = int(_56); +constant int _60 = int(_57); +constant int _61 = _58 ? 2 : 1; +constant int3 _62 = int3(_3, 0, 0); +constant int3 _63 = int3(0, _4, 0); +constant int3 _64 = int3(0, 0, _5); +constant int3 _65 = int3(_62.x, 0, _62.z); +constant int3 _66 = int3(0, _63.y, _63.x); +constant int3 _67 = int3(_64.z, 0, _64.z); +constant int3 _68 = int3(_65.y, _65.x, _66.y); +constant int3 _69 = int3(_67.z, _68.y, _68.z); +constant int _70 = _69.x; +constant int _71 = _69.y; +constant int _72 = _69.z; +constant int _73 = (_70 - _71); +constant int _74 = (_73 * _72); + +constant spvUnsafeArray _33 = spvUnsafeArray({ 0, 0, 0 }); +constant spvUnsafeArray, 3> _34 = spvUnsafeArray, 3>({ spvUnsafeArray({ 0, 0, 0 }), spvUnsafeArray({ 0, 0, 0 }), spvUnsafeArray({ 0, 0, 0 }) }); + +constant int3 _32 = {}; + +kernel void main0(device _7& _8 [[buffer(0)]], device _7& _9 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) +{ + _9._m0[gl_GlobalInvocationID.x] = _8._m0[gl_GlobalInvocationID.x] + ((((1 - _59) * _60) * (_61 - 1)) * _74); +} + diff --git a/shaders-msl/asm/comp/op-spec-constant-op-vector-related.asm.comp b/shaders-msl/asm/comp/op-spec-constant-op-vector-related.asm.comp new file mode 100644 index 00000000..65a7eedd --- /dev/null +++ b/shaders-msl/asm/comp/op-spec-constant-op-vector-related.asm.comp @@ -0,0 +1,107 @@ +OpCapability Shader +OpMemoryModel Logical GLSL450 +OpEntryPoint GLCompute %main "main" %id +OpExecutionMode %main LocalSize 1 1 1 +OpName %main "main" +OpName %id "gl_GlobalInvocationID" +OpDecorate %id BuiltIn GlobalInvocationId +OpDecorate %sc_0 SpecId 0 +OpDecorate %sc_1 SpecId 1 +OpDecorate %sc_2 SpecId 2 +OpDecorate %i32arr ArrayStride 4 +OpDecorate %buf BufferBlock +OpDecorate %indata DescriptorSet 0 +OpDecorate %indata Binding 0 +OpDecorate %outdata DescriptorSet 0 +OpDecorate %outdata Binding 1 +OpDecorate %f32arr ArrayStride 4 +OpMemberDecorate %buf 0 Offset 0 +%bool = OpTypeBool +%void = OpTypeVoid +%voidf = OpTypeFunction %void +%u32 = OpTypeInt 32 0 +%i32 = OpTypeInt 32 1 +%f32 = OpTypeFloat 32 +%uvec3 = OpTypeVector %u32 3 +%fvec3 = OpTypeVector %f32 3 +%uvec3ptr = OpTypePointer Input %uvec3 +%i32ptr = OpTypePointer Uniform %i32 +%f32ptr = OpTypePointer Uniform %f32 +%i32arr = OpTypeRuntimeArray %i32 +%f32arr = OpTypeRuntimeArray %f32 +%ivec3 = OpTypeVector %i32 3 +%zero = OpConstant %i32 0 +%one = OpConstant %i32 1 +%two = OpConstant %i32 2 +%three = OpConstant %i32 3 +%iarr3 = OpTypeArray %i32 %three +%imat3 = OpTypeArray %iarr3 %three +%struct = OpTypeStruct %imat3 +%buf = OpTypeStruct %i32arr +%bufptr = OpTypePointer Uniform %buf +%indata = OpVariable %bufptr Uniform +%outdata = OpVariable %bufptr Uniform +%id = OpVariable %uvec3ptr Input +%ivec3_0 = OpConstantComposite %ivec3 %zero %zero %zero +%vec3_undef = OpUndef %ivec3 +%iarr3_0 = OpConstantComposite %iarr3 %zero %zero %zero +%imat3_0 = OpConstantComposite %imat3 %iarr3_0 %iarr3_0 %iarr3_0 +%struct_0 = OpConstantComposite %struct %imat3_0 +%sc_0 = OpSpecConstant %i32 0 +%sc_1 = OpSpecConstant %i32 0 +%sc_2 = OpSpecConstant %i32 0 +%iarr3_a = OpSpecConstantOp %iarr3 CompositeInsert %sc_0 %iarr3_0 0 +%iarr3_b = OpSpecConstantOp %iarr3 CompositeInsert %sc_1 %iarr3_a 1 +%iarr3_c = OpSpecConstantOp %iarr3 CompositeInsert %sc_2 %iarr3_b 2 +%iarr3_d = OpSpecConstantOp %iarr3 CompositeInsert %sc_1 %iarr3_0 0 +%iarr3_e = OpSpecConstantOp %iarr3 CompositeInsert %sc_2 %iarr3_d 1 +%iarr3_f = OpSpecConstantOp %iarr3 CompositeInsert %sc_0 %iarr3_e 2 +%iarr3_g = OpSpecConstantOp %iarr3 CompositeInsert %sc_2 %iarr3_0 0 +%iarr3_h = OpSpecConstantOp %iarr3 CompositeInsert %sc_0 %iarr3_g 1 +%iarr3_i = OpSpecConstantOp %iarr3 CompositeInsert %sc_1 %iarr3_h 2 +%imat3_a = OpSpecConstantOp %imat3 CompositeInsert %iarr3_c %imat3_0 0 +%imat3_b = OpSpecConstantOp %imat3 CompositeInsert %iarr3_f %imat3_a 1 +%imat3_c = OpSpecConstantOp %imat3 CompositeInsert %iarr3_i %imat3_b 2 +%struct_a = OpSpecConstantOp %struct CompositeInsert %imat3_c %struct_0 0 +%struct_b = OpSpecConstantOp %struct CompositeInsert %sc_2 %struct_a 0 1 2 +%comp_0_0 = OpSpecConstantOp %i32 CompositeExtract %struct_a 0 0 0 +%comp_1_0 = OpSpecConstantOp %i32 CompositeExtract %struct_a 0 1 0 +%comp_0_1 = OpSpecConstantOp %i32 CompositeExtract %struct_a 0 0 1 +%comp_2_2 = OpSpecConstantOp %i32 CompositeExtract %struct_a 0 2 2 +%comp_2_0 = OpSpecConstantOp %i32 CompositeExtract %struct_a 0 2 0 +%comp_1_1 = OpSpecConstantOp %i32 CompositeExtract %struct_a 0 1 1 +%cmpres_0 = OpSpecConstantOp %bool IEqual %comp_0_0 %comp_1_0 +%cmpres_1 = OpSpecConstantOp %bool IEqual %comp_0_1 %comp_2_2 +%cmpres_2 = OpSpecConstantOp %bool IEqual %comp_2_0 %comp_1_1 +%mustbe_0 = OpSpecConstantOp %i32 Select %cmpres_0 %one %zero +%mustbe_1 = OpSpecConstantOp %i32 Select %cmpres_1 %one %zero +%mustbe_2 = OpSpecConstantOp %i32 Select %cmpres_2 %two %one +%sc_vec3_0 = OpSpecConstantOp %ivec3 CompositeInsert %sc_0 %ivec3_0 0 +%sc_vec3_1 = OpSpecConstantOp %ivec3 CompositeInsert %sc_1 %ivec3_0 1 +%sc_vec3_2 = OpSpecConstantOp %ivec3 CompositeInsert %sc_2 %ivec3_0 2 +%sc_vec3_0_s = OpSpecConstantOp %ivec3 VectorShuffle %sc_vec3_0 %vec3_undef 0 0xFFFFFFFF 2 +%sc_vec3_1_s = OpSpecConstantOp %ivec3 VectorShuffle %sc_vec3_1 %vec3_undef 0xFFFFFFFF 1 0 +%sc_vec3_2_s = OpSpecConstantOp %ivec3 VectorShuffle %vec3_undef %sc_vec3_2 5 0xFFFFFFFF 5 +%sc_vec3_01 = OpSpecConstantOp %ivec3 VectorShuffle %sc_vec3_0_s %sc_vec3_1_s 1 0 4 +%sc_vec3_012 = OpSpecConstantOp %ivec3 VectorShuffle %sc_vec3_01 %sc_vec3_2_s 5 1 2 +%sc_ext_0 = OpSpecConstantOp %i32 CompositeExtract %sc_vec3_012 0 +%sc_ext_1 = OpSpecConstantOp %i32 CompositeExtract %sc_vec3_012 1 +%sc_ext_2 = OpSpecConstantOp %i32 CompositeExtract %sc_vec3_012 2 +%sc_sub = OpSpecConstantOp %i32 ISub %sc_ext_0 %sc_ext_1 +%sc_factor = OpSpecConstantOp %i32 IMul %sc_sub %sc_ext_2 +%main = OpFunction %void None %voidf +%label = OpLabel +%subf_a = OpISub %i32 %one %mustbe_0 +%subf_b = OpIMul %i32 %subf_a %mustbe_1 +%subf_c = OpISub %i32 %mustbe_2 %one +%factor = OpIMul %i32 %subf_b %subf_c +%sc_final = OpIMul %i32 %factor %sc_factor +%idval = OpLoad %uvec3 %id +%x = OpCompositeExtract %u32 %idval 0 +%inloc = OpAccessChain %i32ptr %indata %zero %x +%inval = OpLoad %i32 %inloc +%final = OpIAdd %i32 %inval %sc_final +%outloc = OpAccessChain %i32ptr %outdata %zero %x + OpStore %outloc %final + OpReturn + OpFunctionEnd diff --git a/spirv_cross_containers.hpp b/spirv_cross_containers.hpp index 506b069c..1b32870e 100644 --- a/spirv_cross_containers.hpp +++ b/spirv_cross_containers.hpp @@ -210,7 +210,8 @@ public: buffer_capacity = N; } - SmallVector(const T *arg_list_begin, const T *arg_list_end) SPIRV_CROSS_NOEXCEPT : SmallVector() + template + SmallVector(const U *arg_list_begin, const U *arg_list_end) SPIRV_CROSS_NOEXCEPT : SmallVector() { auto count = size_t(arg_list_end - arg_list_begin); reserve(count); @@ -219,7 +220,13 @@ public: this->buffer_size = count; } - SmallVector(std::initializer_list init) SPIRV_CROSS_NOEXCEPT : SmallVector(init.begin(), init.end()) + template + SmallVector(std::initializer_list init) SPIRV_CROSS_NOEXCEPT : SmallVector(init.begin(), init.end()) + { + } + + template + SmallVector(const U (&init)[M]) SPIRV_CROSS_NOEXCEPT : SmallVector(init, init + M) { } diff --git a/spirv_glsl.cpp b/spirv_glsl.cpp index cf621904..ddf1f76f 100644 --- a/spirv_glsl.cpp +++ b/spirv_glsl.cpp @@ -4859,6 +4859,9 @@ string CompilerGLSL::to_expression(uint32_t id, bool register_expression_read) } } + if (expression_is_forwarded(id)) + return constant_expression(c); + return to_name(id); } else if (c.is_used_as_lut) @@ -4930,6 +4933,80 @@ string CompilerGLSL::to_expression(uint32_t id, bool register_expression_read) } } +SmallVector CompilerGLSL::get_composite_constant_ids(ConstantID const_id) +{ + if (auto *constant = maybe_get(const_id)) + { + const auto &type = get(constant->constant_type); + if (is_array(type) || type.basetype == SPIRType::Struct) + return constant->subconstants; + if (is_matrix(type)) + return constant->m.id; + if (is_vector(type)) + return constant->m.c[0].id; + SPIRV_CROSS_THROW("Unexpected scalar constant!"); + } + if (!const_composite_insert_ids.count(const_id)) + SPIRV_CROSS_THROW("Unimplemented for this OpSpecConstantOp!"); + return const_composite_insert_ids[const_id]; +} + +void CompilerGLSL::fill_composite_constant(SPIRConstant &constant, TypeID type_id, + const SmallVector &initializers) +{ + auto &type = get(type_id); + constant.specialization = true; + if (is_array(type) || type.basetype == SPIRType::Struct) + { + constant.subconstants = initializers; + } + else if (is_matrix(type)) + { + constant.m.columns = type.columns; + for (uint32_t i = 0; i < type.columns; ++i) + { + constant.m.id[i] = initializers[i]; + constant.m.c[i].vecsize = type.vecsize; + } + } + else if (is_vector(type)) + { + constant.m.c[0].vecsize = type.vecsize; + for (uint32_t i = 0; i < type.vecsize; ++i) + constant.m.c[0].id[i] = initializers[i]; + } + else + SPIRV_CROSS_THROW("Unexpected scalar in SpecConstantOp CompositeInsert!"); +} + +void CompilerGLSL::set_composite_constant(ConstantID const_id, TypeID type_id, + const SmallVector &initializers) +{ + if (maybe_get(const_id)) + { + const_composite_insert_ids[const_id] = initializers; + return; + } + + auto &constant = set(const_id, type_id); + fill_composite_constant(constant, type_id, initializers); + forwarded_temporaries.insert(const_id); +} + +TypeID CompilerGLSL::get_composite_member_type(TypeID type_id, uint32_t member_idx) +{ + auto &type = get(type_id); + if (is_array(type)) + return type.parent_type; + if (type.basetype == SPIRType::Struct) + return type.member_types[member_idx]; + if (is_matrix(type)) + return type.parent_type; + if (is_vector(type)) + return type.parent_type; + SPIRV_CROSS_THROW("Shouldn't reach lower than vector handling OpSpecConstantOp CompositeInsert!"); +} + string CompilerGLSL::constant_op_expression(const SPIRConstantOp &cop) { auto &type = get(cop.basetype); @@ -5034,10 +5111,21 @@ string CompilerGLSL::constant_op_expression(const SPIRConstantOp &cop) for (uint32_t i = 2; i < uint32_t(cop.arguments.size()); i++) { uint32_t index = cop.arguments[i]; - if (index >= left_components) + if (index == 0xFFFFFFFF) + { + SPIRConstant c; + c.constant_type = type.parent_type; + assert(type.parent_type != ID(0)); + expr += constant_expression(c); + } + else if (index >= left_components) + { expr += right_arg + "." + "xyzw"[index - left_components]; + } else + { expr += left_arg + "." + "xyzw"[index]; + } if (i + 1 < uint32_t(cop.arguments.size())) expr += ", "; @@ -5055,7 +5143,30 @@ string CompilerGLSL::constant_op_expression(const SPIRConstantOp &cop) } case OpCompositeInsert: - SPIRV_CROSS_THROW("OpCompositeInsert spec constant op is not supported."); + { + SmallVector new_init = get_composite_constant_ids(cop.arguments[1]); + uint32_t idx; + uint32_t target_id = cop.self; + uint32_t target_type_id = cop.basetype; + // We have to drill down to the part we want to modify, and create new + // constants for each containing part. + for (idx = 2; idx < cop.arguments.size() - 1; ++idx) + { + uint32_t new_const = ir.increase_bound_by(1); + uint32_t old_const = new_init[cop.arguments[idx]]; + new_init[cop.arguments[idx]] = new_const; + set_composite_constant(target_id, target_type_id, new_init); + new_init = get_composite_constant_ids(old_const); + target_id = new_const; + target_type_id = get_composite_member_type(target_type_id, cop.arguments[idx]); + } + // Now replace the initializer with the one from this instruction. + new_init[cop.arguments[idx]] = cop.arguments[0]; + set_composite_constant(target_id, target_type_id, new_init); + SPIRConstant tmp_const(cop.basetype); + fill_composite_constant(tmp_const, cop.basetype, const_composite_insert_ids[cop.self]); + return constant_expression(tmp_const); + } default: // Some opcodes are unimplemented here, these are currently not possible to test from glslang. @@ -5206,20 +5317,27 @@ string CompilerGLSL::constant_expression(const SPIRConstant &c, bool inside_bloc uint32_t subconstant_index = 0; for (auto &elem : c.subconstants) { - auto &subc = get(elem); - if (subc.specialization) - res += to_name(elem); + if (auto *op = maybe_get(elem)) + { + res += constant_op_expression(*op); + } else { - if (type.array.empty() && type.basetype == SPIRType::Struct) + auto &subc = get(elem); + if (subc.specialization && !expression_is_forwarded(elem)) + res += to_name(elem); + else { - // When we get down to emitting struct members, override the block-like information. - // For constants, we can freely mix and match block-like state. - inside_block_like_struct_scope = - has_member_decoration(type.self, subconstant_index, DecorationOffset); - } + if (type.array.empty() && type.basetype == SPIRType::Struct) + { + // When we get down to emitting struct members, override the block-like information. + // For constants, we can freely mix and match block-like state. + inside_block_like_struct_scope = + has_member_decoration(type.self, subconstant_index, DecorationOffset); + } - res += constant_expression(subc, inside_block_like_struct_scope); + res += constant_expression(subc, inside_block_like_struct_scope); + } } if (&elem != &c.subconstants.back()) diff --git a/spirv_glsl.hpp b/spirv_glsl.hpp index 15ffac29..f8d17259 100644 --- a/spirv_glsl.hpp +++ b/spirv_glsl.hpp @@ -989,6 +989,12 @@ protected: private: void init(); + + SmallVector get_composite_constant_ids(ConstantID const_id); + void fill_composite_constant(SPIRConstant &constant, TypeID type_id, const SmallVector &initializers); + void set_composite_constant(ConstantID const_id, TypeID type_id, const SmallVector &initializers); + TypeID get_composite_member_type(TypeID type_id, uint32_t member_idx); + std::unordered_map> const_composite_insert_ids; }; } // namespace SPIRV_CROSS_NAMESPACE