Merge pull request #2050 from cdavis5e/op-spec-constant-op-composite-insert

MSL: Implement `CompositeInsert` `OpSpecConstantOp`.

6 files changed, 475 insertions, 14 deletions

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 <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+template<typename T, size_t Num>
+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<spvUnsafeArray<int, 3>, 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<int, 3> _36 = spvUnsafeArray<int, 3>({ _3, 0, 0 });
+constant spvUnsafeArray<int, 3> _37 = spvUnsafeArray<int, 3>({ _3, _4, 0 });
+constant spvUnsafeArray<int, 3> _38 = spvUnsafeArray<int, 3>({ _3, _4, _5 });
+constant spvUnsafeArray<int, 3> _39 = spvUnsafeArray<int, 3>({ _4, 0, 0 });
+constant spvUnsafeArray<int, 3> _40 = spvUnsafeArray<int, 3>({ _4, _5, 0 });
+constant spvUnsafeArray<int, 3> _41 = spvUnsafeArray<int, 3>({ _4, _5, _3 });
+constant spvUnsafeArray<int, 3> _42 = spvUnsafeArray<int, 3>({ _5, 0, 0 });
+constant spvUnsafeArray<int, 3> _43 = spvUnsafeArray<int, 3>({ _5, _3, 0 });
+constant spvUnsafeArray<int, 3> _44 = spvUnsafeArray<int, 3>({ _5, _3, _4 });
+constant spvUnsafeArray<spvUnsafeArray<int, 3>, 3> _45 = spvUnsafeArray<spvUnsafeArray<int, 3>, 3>({ spvUnsafeArray<int, 3>({ _3, _4, _5 }), spvUnsafeArray<int, 3>({ 0, 0, 0 }), spvUnsafeArray<int, 3>({ 0, 0, 0 }) });
+constant spvUnsafeArray<spvUnsafeArray<int, 3>, 3> _46 = spvUnsafeArray<spvUnsafeArray<int, 3>, 3>({ spvUnsafeArray<int, 3>({ _3, _4, _5 }), spvUnsafeArray<int, 3>({ _4, _5, _3 }), spvUnsafeArray<int, 3>({ 0, 0, 0 }) });
+constant spvUnsafeArray<spvUnsafeArray<int, 3>, 3> _47 = spvUnsafeArray<spvUnsafeArray<int, 3>, 3>({ spvUnsafeArray<int, 3>({ _3, _4, _5 }), spvUnsafeArray<int, 3>({ _4, _5, _3 }), spvUnsafeArray<int, 3>({ _5, _3, _4 }) });
+constant _29 _48 = _29{ spvUnsafeArray<spvUnsafeArray<int, 3>, 3>({ spvUnsafeArray<int, 3>({ _3, _4, _5 }), spvUnsafeArray<int, 3>({ _4, _5, _3 }), spvUnsafeArray<int, 3>({ _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<int, 3> _33 = spvUnsafeArray<int, 3>({ 0, 0, 0 });
+constant spvUnsafeArray<spvUnsafeArray<int, 3>, 3> _34 = spvUnsafeArray<spvUnsafeArray<int, 3>, 3>({ spvUnsafeArray<int, 3>({ 0, 0, 0 }), spvUnsafeArray<int, 3>({ 0, 0, 0 }), spvUnsafeArray<int, 3>({ 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 <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+template<typename T, size_t Num>
+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<spvUnsafeArray<int, 3>, 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<int, 3> _36 = spvUnsafeArray<int, 3>({ _3, 0, 0 });
+constant spvUnsafeArray<int, 3> _37 = spvUnsafeArray<int, 3>({ _3, _4, 0 });
+constant spvUnsafeArray<int, 3> _38 = spvUnsafeArray<int, 3>({ _3, _4, _5 });
+constant spvUnsafeArray<int, 3> _39 = spvUnsafeArray<int, 3>({ _4, 0, 0 });
+constant spvUnsafeArray<int, 3> _40 = spvUnsafeArray<int, 3>({ _4, _5, 0 });
+constant spvUnsafeArray<int, 3> _41 = spvUnsafeArray<int, 3>({ _4, _5, _3 });
+constant spvUnsafeArray<int, 3> _42 = spvUnsafeArray<int, 3>({ _5, 0, 0 });
+constant spvUnsafeArray<int, 3> _43 = spvUnsafeArray<int, 3>({ _5, _3, 0 });
+constant spvUnsafeArray<int, 3> _44 = spvUnsafeArray<int, 3>({ _5, _3, _4 });
+constant spvUnsafeArray<spvUnsafeArray<int, 3>, 3> _45 = spvUnsafeArray<spvUnsafeArray<int, 3>, 3>({ spvUnsafeArray<int, 3>({ _3, _4, _5 }), spvUnsafeArray<int, 3>({ 0, 0, 0 }), spvUnsafeArray<int, 3>({ 0, 0, 0 }) });
+constant spvUnsafeArray<spvUnsafeArray<int, 3>, 3> _46 = spvUnsafeArray<spvUnsafeArray<int, 3>, 3>({ spvUnsafeArray<int, 3>({ _3, _4, _5 }), spvUnsafeArray<int, 3>({ _4, _5, _3 }), spvUnsafeArray<int, 3>({ 0, 0, 0 }) });
+constant spvUnsafeArray<spvUnsafeArray<int, 3>, 3> _47 = spvUnsafeArray<spvUnsafeArray<int, 3>, 3>({ spvUnsafeArray<int, 3>({ _3, _4, _5 }), spvUnsafeArray<int, 3>({ _4, _5, _3 }), spvUnsafeArray<int, 3>({ _5, _3, _4 }) });
+constant _29 _48 = _29{ spvUnsafeArray<spvUnsafeArray<int, 3>, 3>({ spvUnsafeArray<int, 3>({ _3, _4, _5 }), spvUnsafeArray<int, 3>({ _4, _5, _3 }), spvUnsafeArray<int, 3>({ _5, _3, _4 }) }) };
+constant _29 _49 = _29{ spvUnsafeArray<spvUnsafeArray<int, 3>, 3>({ spvUnsafeArray<int, 3>({ _3, _4, _5 }), spvUnsafeArray<int, 3>({ _4, _5, _5 }), spvUnsafeArray<int, 3>({ _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<int, 3> _33 = spvUnsafeArray<int, 3>({ 0, 0, 0 });
+constant spvUnsafeArray<spvUnsafeArray<int, 3>, 3> _34 = spvUnsafeArray<spvUnsafeArray<int, 3>, 3>({ spvUnsafeArray<int, 3>({ 0, 0, 0 }), spvUnsafeArray<int, 3>({ 0, 0, 0 }), spvUnsafeArray<int, 3>({ 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 <typename U>
+	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<T> init) SPIRV_CROSS_NOEXCEPT : SmallVector(init.begin(), init.end())
+	template <typename U>
+	SmallVector(std::initializer_list<U> init) SPIRV_CROSS_NOEXCEPT : SmallVector(init.begin(), init.end())
+	{
+	}
+
+	template <typename U, size_t M>
+	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<ConstantID> CompilerGLSL::get_composite_constant_ids(ConstantID const_id)
+{
+	if (auto *constant = maybe_get<SPIRConstant>(const_id))
+	{
+		const auto &type = get<SPIRType>(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<ConstantID> &initializers)
+{
+	auto &type = get<SPIRType>(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<ConstantID> &initializers)
+{
+	if (maybe_get<SPIRConstantOp>(const_id))
+	{
+		const_composite_insert_ids[const_id] = initializers;
+		return;
+	}
+
+	auto &constant = set<SPIRConstant>(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<SPIRType>(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<SPIRType>(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<ConstantID> 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<SPIRConstant>(elem);
-			if (subc.specialization)
-				res += to_name(elem);
+			if (auto *op = maybe_get<SPIRConstantOp>(elem))
+			{
+				res += constant_op_expression(*op);
+			}
 			else
 			{
-				if (type.array.empty() && type.basetype == SPIRType::Struct)
+				auto &subc = get<SPIRConstant>(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<ConstantID> get_composite_constant_ids(ConstantID const_id);
+	void fill_composite_constant(SPIRConstant &constant, TypeID type_id, const SmallVector<ConstantID> &initializers);
+	void set_composite_constant(ConstantID const_id, TypeID type_id, const SmallVector<ConstantID> &initializers);
+	TypeID get_composite_member_type(TypeID type_id, uint32_t member_idx);
+	std::unordered_map<uint32_t, SmallVector<ConstantID>> const_composite_insert_ids;
 };
 } // namespace SPIRV_CROSS_NAMESPACE