Merge branch 'master' of github.com:bkaradzic/bgfx

3rdparty/dear-imgui/imgui.h

@@ -2082,26 +2082,26 @@ struct ImFontAtlas
 // ImFontAtlas automatically loads a default embedded font for you when you call GetTexDataAsAlpha8() or GetTexDataAsRGBA32().
 struct ImFont
 {
-    // Members: Hot ~24/32 bytes (for CalcTextSize)
-    ImVector<float>             IndexAdvanceX;      // 12/16 // out //            // Sparse. Glyphs->AdvanceX in a directly indexable way (cache-friendly for CalcTextSize functions which only this this info, and are often bottleneck in large UI).
-    float                       FontSize;           // 4     // in  // <user set> // Height of characters, set during loading (don't change after loading)
+    // Members: Hot ~20/24 bytes (for CalcTextSize)
+    ImVector<float>             IndexAdvanceX;      // 12-16 // out //            // Sparse. Glyphs->AdvanceX in a directly indexable way (cache-friendly for CalcTextSize functions which only this this info, and are often bottleneck in large UI).
     float                       FallbackAdvanceX;   // 4     // out // = FallbackGlyph->AdvanceX
-    ImWchar                     FallbackChar;       // 2     // in  // = '?'      // Replacement glyph if one isn't found. Only set via SetFallbackChar()
+    float                       FontSize;           // 4     // in  //            // Height of characters/line, set during loading (don't change after loading)
 
     // Members: Hot ~36/48 bytes (for CalcTextSize + render loop)
     ImVector<ImWchar>           IndexLookup;        // 12-16 // out //            // Sparse. Index glyphs by Unicode code-point.
     ImVector<ImFontGlyph>       Glyphs;             // 12-16 // out //            // All glyphs.
-    ImVec2                      DisplayOffset;      // 8     // in  // = (0,0)    // Offset font rendering by xx pixels
     const ImFontGlyph*          FallbackGlyph;      // 4-8   // out // = FindGlyph(FontFallbackChar)
+    ImVec2                      DisplayOffset;      // 8     // in  // = (0,0)    // Offset font rendering by xx pixels
 
-    // Members: Cold ~28/40 bytes
+    // Members: Cold ~32/40 bytes
     ImFontAtlas*                ContainerAtlas;     // 4-8   // out //            // What we has been loaded into
     const ImFontConfig*         ConfigData;         // 4-8   // in  //            // Pointer within ContainerAtlas->ConfigData
     short                       ConfigDataCount;    // 2     // in  // ~ 1        // Number of ImFontConfig involved in creating this font. Bigger than 1 when merging multiple font sources into one ImFont.
-    bool                        DirtyLookupTables;  // 1     // out //
+    ImWchar                     FallbackChar;       // 2     // in  // = '?'      // Replacement glyph if one isn't found. Only set via SetFallbackChar()
     float                       Scale;              // 4     // in  // = 1.f      // Base font scale, multiplied by the per-window font scale which you can adjust with SetWindowFontScale()
-    float                       Ascent, Descent;    // 8     // out //            // Ascent: distance from top to bottom of e.g. 'A' [0..FontSize]
+    float                       Ascent, Descent;    // 4+4   // out //            // Ascent: distance from top to bottom of e.g. 'A' [0..FontSize]
     int                         MetricsTotalSurface;// 4     // out //            // Total surface in pixels to get an idea of the font rasterization/texture cost (not exact, we approximate the cost of padding between glyphs)
+    bool                        DirtyLookupTables;  // 1     // out //
 
     // Methods
     IMGUI_API ImFont();

3rdparty/dear-imgui/imgui_widgets.cpp

@@ -730,7 +730,7 @@ void ImGui::Scrollbar(ImGuiLayoutType direction)
     const bool horizontal = (direction == ImGuiLayoutType_Horizontal);
     const ImGuiStyle& style = g.Style;
     const ImGuiID id = GetScrollbarID(direction);
-    
+
     // Render background
     bool other_scrollbar = (horizontal ? window->ScrollbarY : window->ScrollbarX);
     float other_scrollbar_size_w = other_scrollbar ? style.ScrollbarSize : 0.0f;
@@ -2631,6 +2631,8 @@ int ImParseFormatPrecision(const char* fmt, int default_precision)
 // FIXME: Facilitate using this in variety of other situations.
 bool ImGui::InputScalarAsWidgetReplacement(const ImRect& bb, ImGuiID id, const char* label, ImGuiDataType data_type, void* data_ptr, const char* format)
 {
+    IM_UNUSED(id);
+
     ImGuiContext& g = *GImGui;
 
     // On the first frame, g.ScalarAsInputTextId == 0, then on subsequent frames it becomes == id.

3rdparty/glslang/StandAlone/StandAlone.cpp

@@ -534,6 +534,8 @@ void ProcessArguments(std::vector<std::unique_ptr<glslang::TWorkItem>>& workItem
                         ReflectOptions |= EShReflectionSeparateBuffers;
                     } else if (lowerword == "reflect-all-block-variables") {
                         ReflectOptions |= EShReflectionAllBlockVariables;
+                    } else if (lowerword == "reflect-unwrap-io-blocks") {
+                        ReflectOptions |= EShReflectionUnwrapIOBlocks;
                     } else if (lowerword == "resource-set-bindings" ||  // synonyms
                                lowerword == "resource-set-binding"  ||
                                lowerword == "rsb") {
@@ -1538,6 +1540,8 @@ void usage()
            "                                    separately to uniforms\n"
            "  --reflect-all-block-variables     reflect all variables in blocks, whether\n"
            "                                    inactive or active\n"
+           "  --reflect-unwrap-io-blocks        unwrap input/output blocks the same as\n"
+           "                                    uniform blocks\n"
            "  --resource-set-binding [stage] name set binding\n"
            "                                    set descriptor set and binding for\n"
            "                                    individual resources\n"

3rdparty/glslang/Test/baseResults/cppBad.vert.out

@@ -1,9 +1,10 @@
 cppBad.vert
-ERROR: 0:2: 'preprocessor evaluation' : bad expression 
-ERROR: 0:2: '#if' : unexpected tokens following directive 
-ERROR: 0:5: 'string' : End of line in string 
-ERROR: 0:5: '""' : string literals not supported 
-ERROR: 0:5: '' :  syntax error, unexpected INT, expecting COMMA or SEMICOLON
+WARNING: 0:1: '#define' : missing space after macro name 
+ERROR: 0:3: 'preprocessor evaluation' : bad expression 
+ERROR: 0:3: '#if' : unexpected tokens following directive 
+ERROR: 0:6: 'string' : End of line in string 
+ERROR: 0:6: '""' : string literals not supported 
+ERROR: 0:6: '' :  syntax error, unexpected INT, expecting COMMA or SEMICOLON
 ERROR: 5 compilation errors.  No code generated.
 
 

3rdparty/glslang/Test/baseResults/cppBad3.vert.out

@@ -0,0 +1,18 @@
+cppBad3.vert
+ERROR: 0:3: 'macro expansion' : End of input in macro y
+ERROR: 1 compilation errors.  No code generated.
+
+
+Shader version: 100
+ERROR: node is still EOpNull!
+0:?   Linker Objects
+
+
+Linked vertex stage:
+
+ERROR: Linking vertex stage: Missing entry point: Each stage requires one entry point
+
+Shader version: 100
+ERROR: node is still EOpNull!
+0:?   Linker Objects
+

3rdparty/glslang/Test/baseResults/reflection.linked.options.out

@@ -14,8 +14,8 @@ Buffer variable reflection:
 Buffer block reflection:
 
 Pipeline input reflection:
-vertin: offset 0, type 1406, size 0, index 0, binding -1, stages 1
+vertin: offset 0, type 1406, size 1, index 0, binding -1, stages 1
 
 Pipeline output reflection:
-fragout: offset 0, type 1406, size 0, index 0, binding -1, stages 16
+fragout: offset 0, type 1406, size 1, index 0, binding -1, stages 16
 

3rdparty/glslang/Test/baseResults/reflection.options.frag.out

@@ -8,7 +8,7 @@ Buffer variable reflection:
 Buffer block reflection:
 
 Pipeline input reflection:
-inval: offset 0, type 1406, size 0, index 0, binding -1, stages 16
+inval: offset 0, type 1406, size 1, index 0, binding -1, stages 16
 
 Pipeline output reflection:
 

3rdparty/glslang/Test/baseResults/reflection.options.geom.out

@@ -0,0 +1,25 @@
+reflection.options.geom
+Uniform reflection:
+
+Uniform block reflection:
+
+Buffer variable reflection:
+
+Buffer block reflection:
+
+Pipeline input reflection:
+gl_PerVertex.gl_Position: offset 0, type 8b52, size 1, index 0, binding -1, stages 8
+gl_PerVertex.gl_PointSize: offset 0, type 1406, size 1, index 0, binding -1, stages 8
+gl_PerVertex.gl_ClipDistance[0]: offset 0, type 1406, size 1, index 0, binding -1, stages 8
+block.Color: offset 0, type 8b50, size 1, index 0, binding -1, stages 8
+block.Texcoord: offset 0, type 8b50, size 1, index 0, binding -1, stages 8
+block.in_a: offset 0, type 8b54, size 1, index 0, binding -1, stages 8
+
+Pipeline output reflection:
+gl_Position: offset 0, type 8b52, size 1, index 0, binding -1, stages 8
+gl_PointSize: offset 0, type 1406, size 1, index 0, binding -1, stages 8
+gl_ClipDistance[0]: offset 0, type 1406, size 1, index 0, binding -1, stages 8
+block.Color: offset 0, type 8b52, size 1, index 0, binding -1, stages 8
+block.a: offset 0, type 8b52, size 1, index 0, binding -1, stages 8
+block.b[0]: offset 0, type 8b50, size 3, index 0, binding -1, stages 8
+

3rdparty/glslang/Test/baseResults/reflection.options.vert.out

@@ -33,8 +33,11 @@ VertexCollection: offset -1, type ffffffff, size 400, index -1, binding -1, stag
 MultipleArrays: offset -1, type ffffffff, size 500, index -1, binding -1, stages 1, numMembers 9
 
 Pipeline input reflection:
-gl_InstanceID: offset 0, type 1404, size 0, index 0, binding -1, stages 1
+gl_InstanceID: offset 0, type 1404, size 1, index 0, binding -1, stages 1
 
 Pipeline output reflection:
-outval: offset 0, type 1406, size 0, index 0, binding -1, stages 1
+outval.val: offset 0, type 1406, size 1, index 0, binding -1, stages 1
+outval.a: offset 0, type 8b51, size 1, index 0, binding -1, stages 1
+outval.b[0]: offset 0, type 8b50, size 4, index 0, binding -1, stages 1
+outval.c: offset 0, type 8b5a, size 1, index 0, binding -1, stages 1
 

3rdparty/glslang/Test/baseResults/spv.meshShaderBuiltins.mesh.out

@@ -1,7 +1,7 @@
 spv.meshShaderBuiltins.mesh
 // Module Version 10000
 // Generated by (magic number): 80007
-// Id's are bound by 146
+// Id's are bound by 148
 
                               Capability ClipDistance
                               Capability CullDistance
@@ -14,7 +14,7 @@ spv.meshShaderBuiltins.mesh
                               Extension  "SPV_NV_viewport_array2"
                1:             ExtInstImport  "GLSL.std.450"
                               MemoryModel Logical GLSL450
-                              EntryPoint MeshNV 4  "main" 11 17 34 88 128 139 143
+                              EntryPoint MeshNV 4  "main" 11 17 34 88 129 142 146
                               ExecutionMode 4 LocalSize 32 1 1
                               ExecutionMode 4 OutputVertices 81
                               ExecutionMode 4 OutputPrimitivesNV 32
@@ -43,9 +43,9 @@ spv.meshShaderBuiltins.mesh
                               MemberName 84(gl_MeshPerPrimitiveNV) 4  "gl_LayerPerViewNV"
                               MemberName 84(gl_MeshPerPrimitiveNV) 5  "gl_ViewportMaskPerViewNV"
                               Name 88  "gl_MeshPrimitivesNV"
-                              Name 128  "gl_PrimitiveIndicesNV"
-                              Name 139  "gl_DrawID"
-                              Name 143  "gl_PrimitiveCountNV"
+                              Name 129  "gl_PrimitiveIndicesNV"
+                              Name 142  "gl_DrawID"
+                              Name 146  "gl_PrimitiveCountNV"
                               Decorate 11(gl_LocalInvocationID) BuiltIn LocalInvocationId
                               Decorate 17(gl_WorkGroupID) BuiltIn WorkgroupId
                               MemberDecorate 30(gl_MeshPerVertexNV) 0 BuiltIn Position
@@ -74,10 +74,10 @@ spv.meshShaderBuiltins.mesh
                               MemberDecorate 84(gl_MeshPerPrimitiveNV) 5 PerViewNV
                               MemberDecorate 84(gl_MeshPerPrimitiveNV) 5 BuiltIn ViewportMaskPerViewNV
                               Decorate 84(gl_MeshPerPrimitiveNV) Block
-                              Decorate 128(gl_PrimitiveIndicesNV) BuiltIn PrimitiveIndicesNV
-                              Decorate 139(gl_DrawID) BuiltIn DrawIndex
-                              Decorate 143(gl_PrimitiveCountNV) BuiltIn PrimitiveCountNV
-                              Decorate 145 BuiltIn WorkgroupSize
+                              Decorate 129(gl_PrimitiveIndicesNV) BuiltIn PrimitiveIndicesNV
+                              Decorate 142(gl_DrawID) BuiltIn DrawIndex
+                              Decorate 146(gl_PrimitiveCountNV) BuiltIn PrimitiveCountNV
+                              Decorate 147 BuiltIn WorkgroupSize
                2:             TypeVoid
                3:             TypeFunction 2
                6:             TypeInt 32 0
@@ -130,17 +130,18 @@ spv.meshShaderBuiltins.mesh
               94:     36(int) Constant 7
               97:     36(int) Constant 8
              100:     36(int) Constant 9
-             126:             TypeArray 6(int) 31
-             127:             TypePointer Output 126
-128(gl_PrimitiveIndicesNV):    127(ptr) Variable Output
-             129:      6(int) Constant 257
-             130:             TypePointer Output 6(int)
-             138:             TypePointer Input 36(int)
-  139(gl_DrawID):    138(ptr) Variable Input
-             142:      6(int) Constant 16909060
-143(gl_PrimitiveCountNV):    130(ptr) Variable Output
-             144:      6(int) Constant 96
-             145:    9(ivec3) ConstantComposite 85 27 27
+             126:      6(int) Constant 96
+             127:             TypeArray 6(int) 126
+             128:             TypePointer Output 127
+129(gl_PrimitiveIndicesNV):    128(ptr) Variable Output
+             130:      6(int) Constant 257
+             131:             TypePointer Output 6(int)
+             133:     36(int) Constant 95
+             141:             TypePointer Input 36(int)
+  142(gl_DrawID):    141(ptr) Variable Input
+             145:      6(int) Constant 16909060
+146(gl_PrimitiveCountNV):    131(ptr) Variable Output
+             147:    9(ivec3) ConstantComposite 85 27 27
          4(main):           2 Function None 3
                5:             Label
           8(iid):      7(ptr) Variable Function
@@ -239,19 +240,21 @@ spv.meshShaderBuiltins.mesh
                               Store 125 124
                               MemoryBarrier 27 55
                               ControlBarrier 56 56 55
-             131:    130(ptr) AccessChain 128(gl_PrimitiveIndicesNV) 37
-                              Store 131 129
-             132:      6(int) Load 16(gid)
-             133:      6(int) Load 16(gid)
-             134:      6(int) ISub 133 27
-             135:    130(ptr) AccessChain 128(gl_PrimitiveIndicesNV) 134
-             136:      6(int) Load 135
-             137:    130(ptr) AccessChain 128(gl_PrimitiveIndicesNV) 132
-                              Store 137 136
-             140:     36(int) Load 139(gl_DrawID)
-             141:      6(int) Bitcast 140
-             142:         141 WritePackedPrimitiveIndices4x8NV
-                              Store 143(gl_PrimitiveCountNV) 144
+             132:    131(ptr) AccessChain 129(gl_PrimitiveIndicesNV) 37
+                              Store 132 130
+             134:    131(ptr) AccessChain 129(gl_PrimitiveIndicesNV) 133
+                              Store 134 56
+             135:      6(int) Load 16(gid)
+             136:      6(int) Load 16(gid)
+             137:      6(int) ISub 136 27
+             138:    131(ptr) AccessChain 129(gl_PrimitiveIndicesNV) 137
+             139:      6(int) Load 138
+             140:    131(ptr) AccessChain 129(gl_PrimitiveIndicesNV) 135
+                              Store 140 139
+             143:     36(int) Load 142(gl_DrawID)
+             144:      6(int) Bitcast 143
+             145:         144 WritePackedPrimitiveIndices4x8NV
+                              Store 146(gl_PrimitiveCountNV) 126
                               MemoryBarrier 27 55
                               ControlBarrier 56 56 55
                               Return

3rdparty/glslang/Test/baseResults/spv.meshShaderRedeclBuiltins.mesh.out

@@ -1,7 +1,7 @@
 spv.meshShaderRedeclBuiltins.mesh
 // Module Version 10000
 // Generated by (magic number): 80007
-// Id's are bound by 120
+// Id's are bound by 129
 
                               Capability ClipDistance
                               Capability CullDistance
@@ -12,7 +12,7 @@ spv.meshShaderRedeclBuiltins.mesh
                               Extension  "SPV_NV_viewport_array2"
                1:             ExtInstImport  "GLSL.std.450"
                               MemoryModel Logical GLSL450
-                              EntryPoint MeshNV 4  "main" 11 17 28 81
+                              EntryPoint MeshNV 4  "main" 11 17 28 81 122 127
                               ExecutionMode 4 LocalSize 32 1 1
                               ExecutionMode 4 OutputVertices 81
                               ExecutionMode 4 OutputPrimitivesNV 32
@@ -36,6 +36,8 @@ spv.meshShaderRedeclBuiltins.mesh
                               MemberName 77(gl_MeshPerPrimitiveNV) 2  "gl_ViewportIndex"
                               MemberName 77(gl_MeshPerPrimitiveNV) 3  "gl_ViewportMask"
                               Name 81  "gl_MeshPrimitivesNV"
+                              Name 122  "gl_PrimitiveIndicesNV"
+                              Name 127  "gl_PrimitiveCountNV"
                               Decorate 11(gl_LocalInvocationID) BuiltIn LocalInvocationId
                               Decorate 17(gl_WorkGroupID) BuiltIn WorkgroupId
                               MemberDecorate 24(gl_MeshPerVertexNV) 0 BuiltIn Position
@@ -52,7 +54,9 @@ spv.meshShaderRedeclBuiltins.mesh
                               MemberDecorate 77(gl_MeshPerPrimitiveNV) 3 PerPrimitiveNV
                               MemberDecorate 77(gl_MeshPerPrimitiveNV) 3 BuiltIn ViewportMaskNV
                               Decorate 77(gl_MeshPerPrimitiveNV) Block
-                              Decorate 119 BuiltIn WorkgroupSize
+                              Decorate 122(gl_PrimitiveIndicesNV) BuiltIn PrimitiveIndicesNV
+                              Decorate 127(gl_PrimitiveCountNV) BuiltIn PrimitiveCountNV
+                              Decorate 128 BuiltIn WorkgroupSize
                2:             TypeVoid
                3:             TypeFunction 2
                6:             TypeInt 32 0
@@ -98,7 +102,14 @@ spv.meshShaderRedeclBuiltins.mesh
               87:     30(int) Constant 7
               90:     30(int) Constant 8
               93:     30(int) Constant 9
-             119:    9(ivec3) ConstantComposite 78 49 49
+             119:      6(int) Constant 96
+             120:             TypeArray 6(int) 119
+             121:             TypePointer Output 120
+122(gl_PrimitiveIndicesNV):    121(ptr) Variable Output
+             123:             TypePointer Output 6(int)
+             125:     30(int) Constant 95
+127(gl_PrimitiveCountNV):    123(ptr) Variable Output
+             128:    9(ivec3) ConstantComposite 78 49 49
          4(main):           2 Function None 3
                5:             Label
           8(iid):      7(ptr) Variable Function
@@ -197,5 +208,10 @@ spv.meshShaderRedeclBuiltins.mesh
                               Store 118 117
                               MemoryBarrier 49 50
                               ControlBarrier 51 51 50
+             124:    123(ptr) AccessChain 122(gl_PrimitiveIndicesNV) 31
+                              Store 124 49
+             126:    123(ptr) AccessChain 122(gl_PrimitiveIndicesNV) 125
+                              Store 126 51
+                              Store 127(gl_PrimitiveCountNV) 119
                               Return
                               FunctionEnd

3rdparty/glslang/Test/cppBad.vert

@@ -1,4 +1,5 @@
-#define m#0#
+#define n#0#
+#define m #0#
 #if m
 #endif
 #define n()

3rdparty/glslang/Test/cppBad3.vert

@@ -0,0 +1,3 @@
+#define f =y(.
+#define y(m)
+y(f)

3rdparty/glslang/Test/reflection.options.geom

@@ -0,0 +1,31 @@
+#version 330 core
+
+precision highp float;
+
+layout(triangles) in;
+layout(triangle_strip, max_vertices = 4) out;
+
+in block
+{
+    vec2 Color;
+    vec2 Texcoord;
+    flat ivec3 in_a;
+} In[];
+
+out block
+{
+    vec4 Color;
+    vec4 a;
+    vec2 b[3];
+} Out;
+
+void main()
+{
+    for(int i = 0; i < gl_in.length(); ++i)
+    {
+        gl_Position = gl_in[i].gl_Position;
+        Out.Color = vec4(In[i].Color, 0, 1);
+        EmitVertex();
+    }
+    EndPrimitive();
+}

3rdparty/glslang/Test/reflection.options.vert

@@ -26,7 +26,14 @@ uniform UBO {
     uvec4 unused;
 } ubo;
 
-out float outval;
+struct OutputStruct {
+    float val;
+    vec3 a;
+    vec2 b[4];
+    mat2x2 c;
+};
+
+out OutputStruct outval;
 
 void main()
 {
@@ -40,5 +47,5 @@ void main()
     f += ubo.verts[gl_InstanceID].position[0];
     f += ubo.flt[gl_InstanceID];
     TriangleInfo tlocal[5] = t;
-    outval = f;
+    outval.val = f;
 }

3rdparty/glslang/Test/runtests

@@ -32,15 +32,17 @@ diff -b $BASEDIR/badMacroArgs.frag.out $TARGETDIR/badMacroArgs.frag.out || HASER
 echo Running reflection...
 $EXE -l -q -C reflection.vert > $TARGETDIR/reflection.vert.out
 diff -b $BASEDIR/reflection.vert.out $TARGETDIR/reflection.vert.out || HASERROR=1
-$EXE -l -q -C --reflect-strict-array-suffix --reflect-basic-array-suffix --reflect-intermediate-io --reflect-separate-buffers --reflect-all-block-variables reflection.options.vert > $TARGETDIR/reflection.options.vert.out
+$EXE -l -q -C --reflect-strict-array-suffix --reflect-basic-array-suffix --reflect-intermediate-io --reflect-separate-buffers --reflect-all-block-variables --reflect-unwrap-io-blocks reflection.options.vert > $TARGETDIR/reflection.options.vert.out
 diff -b $BASEDIR/reflection.options.vert.out $TARGETDIR/reflection.options.vert.out || HASERROR=1
 $EXE -l -q -C reflection.frag > $TARGETDIR/reflection.frag.out
 diff -b $BASEDIR/reflection.frag.out $TARGETDIR/reflection.frag.out || HASERROR=1
-$EXE -l -q -C --reflect-strict-array-suffix --reflect-basic-array-suffix --reflect-intermediate-io --reflect-separate-buffers --reflect-all-block-variables reflection.frag > $TARGETDIR/reflection.options.frag.out
+$EXE -l -q -C --reflect-strict-array-suffix --reflect-basic-array-suffix --reflect-intermediate-io --reflect-separate-buffers --reflect-all-block-variables --reflect-unwrap-io-blocks reflection.frag > $TARGETDIR/reflection.options.frag.out
 diff -b $BASEDIR/reflection.options.frag.out $TARGETDIR/reflection.options.frag.out || HASERROR=1
+$EXE -l -q -C --reflect-strict-array-suffix --reflect-basic-array-suffix --reflect-intermediate-io --reflect-separate-buffers --reflect-all-block-variables --reflect-unwrap-io-blocks reflection.options.geom > $TARGETDIR/reflection.options.geom.out
+diff -b $BASEDIR/reflection.options.geom.out $TARGETDIR/reflection.options.geom.out || HASERROR=1
 $EXE -l -q -C reflection.linked.vert reflection.linked.frag > $TARGETDIR/reflection.linked.out
 diff -b $BASEDIR/reflection.linked.out $TARGETDIR/reflection.linked.out || HASERROR=1
-$EXE -l -q -C --reflect-strict-array-suffix --reflect-basic-array-suffix --reflect-intermediate-io --reflect-separate-buffers --reflect-all-block-variables reflection.linked.vert reflection.linked.frag > $TARGETDIR/reflection.linked.options.out
+$EXE -l -q -C --reflect-strict-array-suffix --reflect-basic-array-suffix --reflect-intermediate-io --reflect-separate-buffers --reflect-all-block-variables --reflect-unwrap-io-blocks reflection.linked.vert reflection.linked.frag > $TARGETDIR/reflection.linked.options.out
 diff -b $BASEDIR/reflection.linked.options.out $TARGETDIR/reflection.linked.options.out || HASERROR=1
 $EXE -D -Od -e flizv -l -q -C -V -Od hlsl.reflection.vert > $TARGETDIR/hlsl.reflection.vert.out
 diff -b $BASEDIR/hlsl.reflection.vert.out $TARGETDIR/hlsl.reflection.vert.out || HASERROR=1

3rdparty/glslang/Test/spv.meshShaderBuiltins.mesh

@@ -50,8 +50,9 @@ void main()
 
     BARRIER();
 
-    // should truncate 257 -> 1
-    gl_PrimitiveIndicesNV[0] = 257;
+    // check bound limits
+    gl_PrimitiveIndicesNV[0] = 257; // should truncate 257 -> 1
+    gl_PrimitiveIndicesNV[(MAX_PRIM * 3) - 1] = 2;
     gl_PrimitiveIndicesNV[gid] = gl_PrimitiveIndicesNV[gid-1];
 
     // writes 4 indices at offset gl_DrawID

3rdparty/glslang/Test/spv.meshShaderRedeclBuiltins.mesh

@@ -22,14 +22,16 @@ out gl_MeshPerVertexNV {
     float gl_PointSize;
     float gl_ClipDistance[4];
     float gl_CullDistance[4];
-} gl_MeshVerticesNV[];
+} gl_MeshVerticesNV[MAX_VER];                   // explicitly sized to MAX_VER
 
 perprimitiveNV out gl_MeshPerPrimitiveNV {
     int gl_PrimitiveID;
     int gl_Layer;
     int gl_ViewportIndex;
     int gl_ViewportMask[];
-} gl_MeshPrimitivesNV[];
+} gl_MeshPrimitivesNV[];                        // implicitly sized to MAX_PRIM
+
+out uint gl_PrimitiveIndicesNV[MAX_PRIM*3];     // explicitly sized to MAX_PRIM * 3
 
 void main()
 {
@@ -63,4 +65,9 @@ void main()
     gl_MeshPrimitivesNV[iid+1].gl_ViewportMask[0] = gl_MeshPrimitivesNV[iid].gl_ViewportMask[0];
 
     BARRIER();
+
+    // check bound limits
+    gl_PrimitiveIndicesNV[0] = 1;
+    gl_PrimitiveIndicesNV[(MAX_PRIM * 3) - 1] = 2;
+    gl_PrimitiveCountNV = MAX_PRIM * 3;
 }

3rdparty/glslang/glslang/MachineIndependent/ParseHelper.cpp

@@ -572,7 +572,7 @@ void TParseContext::handleIoResizeArrayAccess(const TSourceLoc& /*loc*/, TInterm
 
     // fix array size, if it can be fixed and needs to be fixed (will allow variable indexing)
     if (symbolNode->getType().isUnsizedArray()) {
-        int newSize = getIoArrayImplicitSize(symbolNode->getType().getQualifier().isPerPrimitive());
+        int newSize = getIoArrayImplicitSize(symbolNode->getType().getQualifier());
         if (newSize > 0)
             symbolNode->getWritableType().changeOuterArraySize(newSize);
     }
@@ -586,59 +586,80 @@ void TParseContext::handleIoResizeArrayAccess(const TSourceLoc& /*loc*/, TInterm
 // Types without an array size will be given one.
 // Types already having a size that is wrong will get an error.
 //
-void TParseContext::checkIoArraysConsistency(const TSourceLoc& loc, bool tailOnly, bool isPerPrimitive)
+void TParseContext::checkIoArraysConsistency(const TSourceLoc &loc, bool tailOnly)
 {
-    int requiredSize = getIoArrayImplicitSize(isPerPrimitive);
-    if (requiredSize == 0)
-        return;
+    int requiredSize = 0;
+    TString featureString;
+    size_t listSize = ioArraySymbolResizeList.size();
+    size_t i = 0;
 
-    const char* feature;
-    if (language == EShLangGeometry)
-        feature = TQualifier::getGeometryString(intermediate.getInputPrimitive());
-    else if (language == EShLangTessControl
-#ifdef NV_EXTENSIONS
-          || language == EShLangFragment
-#endif
-        )
+    // If tailOnly = true, only check the last array symbol in the list.
+    if (tailOnly) {
+        i = listSize - 1;
+    }
+    for (bool firstIteration = true; i < listSize; ++i) {
+        TType &type = ioArraySymbolResizeList[i]->getWritableType();
 
-        feature = "vertices";
+        // As I/O array sizes don't change, fetch requiredSize only once,
+        // except for mesh shaders which could have different I/O array sizes based on type qualifiers.
+        if (firstIteration
 #ifdef NV_EXTENSIONS
-     else if (language == EShLangMeshNV) {
-        feature = isPerPrimitive ? "max_primitives" : "max_vertices";
-     }
+            || (language == EShLangMeshNV)
 #endif
-    else
-        feature = "unknown";
+        )
+        {
+            requiredSize = getIoArrayImplicitSize(type.getQualifier(), &featureString);
+            if (requiredSize == 0)
+                break;
+            firstIteration = false;
+        }
 
-    if (tailOnly) {
-        checkIoArrayConsistency(loc, requiredSize, feature, ioArraySymbolResizeList.back()->getWritableType(), ioArraySymbolResizeList.back()->getName());
-        return;
+        checkIoArrayConsistency(loc, requiredSize, featureString.c_str(), type,
+                                ioArraySymbolResizeList[i]->getName());
     }
-
-    for (size_t i = 0; i < ioArraySymbolResizeList.size(); ++i)
-        checkIoArrayConsistency(loc, requiredSize, feature, ioArraySymbolResizeList[i]->getWritableType(), ioArraySymbolResizeList[i]->getName());
 }
 
-int TParseContext::getIoArrayImplicitSize(bool isPerPrimitive) const
+int TParseContext::getIoArrayImplicitSize(const TQualifier &qualifier, TString *featureString) const
 {
-    if (language == EShLangGeometry)
-        return TQualifier::mapGeometryToSize(intermediate.getInputPrimitive());
-    else if (language == EShLangTessControl)
-        return intermediate.getVertices() != TQualifier::layoutNotSet ? intermediate.getVertices() : 0;
+    int expectedSize = 0;
+    TString str = "unknown";
+    unsigned int maxVertices = intermediate.getVertices() != TQualifier::layoutNotSet ? intermediate.getVertices() : 0;
+
+    if (language == EShLangGeometry) {
+        expectedSize = TQualifier::mapGeometryToSize(intermediate.getInputPrimitive());
+        str = TQualifier::getGeometryString(intermediate.getInputPrimitive());
+    }
+    else if (language == EShLangTessControl) {
+        expectedSize = maxVertices;
+        str = "vertices";
+    }
 #ifdef NV_EXTENSIONS
-    else if (language == EShLangFragment)
-        return 3; //Number of vertices for Fragment shader is always three.
+    else if (language == EShLangFragment) {
+        // Number of vertices for Fragment shader is always three.
+        expectedSize = 3;
+        str = "vertices";
+    }
     else if (language == EShLangMeshNV) {
-        if (isPerPrimitive) {
-            return intermediate.getPrimitives() != TQualifier::layoutNotSet ? intermediate.getPrimitives() : 0;
-        } else {
-            return intermediate.getVertices() != TQualifier::layoutNotSet ? intermediate.getVertices() : 0;
+        unsigned int maxPrimitives =
+            intermediate.getPrimitives() != TQualifier::layoutNotSet ? intermediate.getPrimitives() : 0;
+        if (qualifier.builtIn == EbvPrimitiveIndicesNV) {
+            expectedSize = maxPrimitives * TQualifier::mapGeometryToSize(intermediate.getOutputPrimitive());
+            str = "max_primitives*";
+            str += TQualifier::getGeometryString(intermediate.getOutputPrimitive());
+        }
+        else if (qualifier.isPerPrimitive()) {
+            expectedSize = maxPrimitives;
+            str = "max_primitives";
+        }
+        else {
+            expectedSize = maxVertices;
+            str = "max_vertices";
         }
     }
 #endif
-
-    else
-        return 0;
+    if (featureString)
+        *featureString = str;
+    return expectedSize;
 }
 
 void TParseContext::checkIoArrayConsistency(const TSourceLoc& loc, int requiredSize, const char* feature, TType& type, const TString& name)
@@ -1386,7 +1407,7 @@ TIntermTyped* TParseContext::handleLengthMethod(const TSourceLoc& loc, TFunction
 #endif
                         )
                     {
-                        length = getIoArrayImplicitSize(type.getQualifier().isPerPrimitive());
+                        length = getIoArrayImplicitSize(type.getQualifier());
                     }
                 }
                 if (length == 0) {
@@ -3730,7 +3751,7 @@ void TParseContext::declareArray(const TSourceLoc& loc, const TString& identifie
             if (! symbolTable.atBuiltInLevel()) {
                 if (isIoResizeArray(type)) {
                     ioArraySymbolResizeList.push_back(symbol);
-                    checkIoArraysConsistency(loc, true, type.getQualifier().isPerPrimitive());
+                    checkIoArraysConsistency(loc, true);
                 } else
                     fixIoArraySize(loc, symbol->getWritableType());
             }
@@ -3783,7 +3804,7 @@ void TParseContext::declareArray(const TSourceLoc& loc, const TString& identifie
     existingType.updateArraySizes(type);
 
     if (isIoResizeArray(type))
-        checkIoArraysConsistency(loc, false, type.getQualifier().isPerPrimitive());
+        checkIoArraysConsistency(loc);
 }
 
 // Policy and error check for needing a runtime sized array.
@@ -3939,6 +3960,7 @@ TSymbol* TParseContext::redeclareBuiltinVariable(const TSourceLoc& loc, const TS
 #ifdef NV_EXTENSIONS
          identifier == "gl_SampleMask"                                                              ||
          identifier == "gl_Layer"                                                                   ||
+         identifier == "gl_PrimitiveIndicesNV"                                                      ||
 #endif
          identifier == "gl_TexCoord") {
 
@@ -4018,7 +4040,11 @@ TSymbol* TParseContext::redeclareBuiltinVariable(const TSourceLoc& loc, const TS
                     error(loc, "all redeclarations must use the same depth layout on", "redeclaration", symbol->getName().c_str());
             }
         }
-        else if (identifier == "gl_FragStencilRefARB") {
+        else if (
+#ifdef NV_EXTENSIONS
+            identifier == "gl_PrimitiveIndicesNV" ||
+#endif
+            identifier == "gl_FragStencilRefARB") {
             if (qualifier.hasLayout())
                 error(loc, "cannot apply layout qualifier to", "redeclaration", symbol->getName().c_str());
             if (qualifier.storage != EvqVaryingOut)
@@ -4270,7 +4296,7 @@ void TParseContext::redeclareBuiltinBlock(const TSourceLoc& loc, TTypeList& newT
     // Tracking for implicit sizing of array
     if (isIoResizeArray(block->getType())) {
         ioArraySymbolResizeList.push_back(block);
-        checkIoArraysConsistency(loc, true, block->getType().getQualifier().isPerPrimitive());
+        checkIoArraysConsistency(loc, true);
     } else if (block->getType().isArray())
         fixIoArraySize(loc, block->getWritableType());
 
@@ -7102,7 +7128,7 @@ void TParseContext::declareBlock(const TSourceLoc& loc, TTypeList& typeList, con
     // fix up
     if (isIoResizeArray(blockType)) {
         ioArraySymbolResizeList.push_back(&variable);
-        checkIoArraysConsistency(loc, true, blockType.getQualifier().isPerPrimitive());
+        checkIoArraysConsistency(loc, true);
     } else
         fixIoArraySize(loc, variable.getWritableType());
 
@@ -7685,6 +7711,14 @@ void TParseContext::updateStandaloneQualifierDefaults(const TSourceLoc& loc, con
         else
             error(loc, "can only apply to 'in'", "derivative_group_linearNV", "");
     }
+    // Check mesh out array sizes, once all the necessary out qualifiers are defined.
+    if ((language == EShLangMeshNV) &&
+        (intermediate.getVertices() != TQualifier::layoutNotSet) &&
+        (intermediate.getPrimitives() != TQualifier::layoutNotSet) &&
+        (intermediate.getOutputPrimitive() != ElgNone))
+    {
+        checkIoArraysConsistency(loc);
+    }
 #endif 
     const TQualifier& qualifier = publicType.qualifier;
 
@@ -7835,3 +7869,4 @@ TIntermNode* TParseContext::addSwitch(const TSourceLoc& loc, TIntermTyped* expre
 }
 
 } // end namespace glslang
+

3rdparty/glslang/glslang/MachineIndependent/ParseHelper.h

@@ -299,8 +299,8 @@ public:
     void fixIoArraySize(const TSourceLoc&, TType&);
     void ioArrayCheck(const TSourceLoc&, const TType&, const TString& identifier);
     void handleIoResizeArrayAccess(const TSourceLoc&, TIntermTyped* base);
-    void checkIoArraysConsistency(const TSourceLoc&, bool tailOnly = false, bool isPerPrimitive = false);
-    int getIoArrayImplicitSize(bool isPerPrimitive = false) const;
+    void checkIoArraysConsistency(const TSourceLoc&, bool tailOnly = false);
+    int getIoArrayImplicitSize(const TQualifier&, TString* featureString = nullptr) const;
     void checkIoArrayConsistency(const TSourceLoc&, int requiredSize, const char* feature, TType&, const TString&);
 
     TIntermTyped* handleBinaryMath(const TSourceLoc&, const char* str, TOperator op, TIntermTyped* left, TIntermTyped* right);

3rdparty/glslang/glslang/MachineIndependent/Versions.cpp

@@ -1025,7 +1025,7 @@ void TParseVersions::float16OpaqueCheck(const TSourceLoc& loc, const char* op, b
 void TParseVersions::explicitInt16Check(const TSourceLoc& loc, const char* op, bool builtIn)
 {
     if (! builtIn) {
-    	const char* const extensions[] = {
+        const char* const extensions[] = {
 #if AMD_EXTENSIONS
                                            E_GL_AMD_gpu_shader_int16,
 #endif

3rdparty/glslang/glslang/MachineIndependent/preprocessor/Pp.cpp

@@ -147,6 +147,10 @@ int TPpContext::CPPdefine(TPpToken* ppToken)
         }
 
         token = scanToken(ppToken);
+    } else if (token != '\n' && token != EndOfInput && !ppToken->space) {
+        parseContext.ppWarn(ppToken->loc, "missing space after macro name", "#define", "");
+
+        return token;
     }
 
     // record the definition of the macro

3rdparty/glslang/glslang/MachineIndependent/preprocessor/PpScanner.cpp

@@ -1056,7 +1056,7 @@ int TPpContext::tokenize(TPpToken& ppToken)
         // Handle token-pasting logic
         token = tokenPaste(token, ppToken);
 
-        if (token == EndOfInput) {
+        if (token == EndOfInput || token == tMarkerInput::marker) {
             missingEndifCheck();
             return EndOfInput;
         }

3rdparty/glslang/glslang/MachineIndependent/reflection.cpp

@@ -98,46 +98,46 @@ public:
         }
     }
 
-    void addPipeInput(const TIntermSymbol& base)
+    void addPipeIOVariable(const TIntermSymbol& base)
     {
         if (processedDerefs.find(&base) == processedDerefs.end()) {
             processedDerefs.insert(&base);
 
             const TString &name = base.getName();
             const TType &type = base.getType();
+            const bool input = base.getQualifier().isPipeInput();
 
-            TReflection::TNameToIndex::const_iterator it = reflection.nameToIndex.find(name.c_str());
-            if (it == reflection.nameToIndex.end()) {
-                reflection.nameToIndex[name.c_str()] = (int)reflection.indexToPipeInput.size();
-                reflection.indexToPipeInput.push_back(TObjectReflection(name.c_str(), type, 0, mapToGlType(type), 0, 0));
+            TReflection::TMapIndexToReflection &ioItems =
+                input ? reflection.indexToPipeInput : reflection.indexToPipeOutput;
 
-                EShLanguageMask& stages = reflection.indexToPipeInput.back().stages;
-                stages = static_cast<EShLanguageMask>(stages | 1 << intermediate.getStage());
-            } else {
-                EShLanguageMask& stages = reflection.indexToPipeInput[it->second].stages;
-                stages = static_cast<EShLanguageMask>(stages | 1 << intermediate.getStage());
-            }
-        }
-    }
+            if (reflection.options & EShReflectionUnwrapIOBlocks) {
+                bool anonymous = IsAnonymous(name);
 
-    void addPipeOutput(const TIntermSymbol& base)
-    {
-        if (processedDerefs.find(&base) == processedDerefs.end()) {
-            processedDerefs.insert(&base);
-
-            const TString &name = base.getName();
-            const TType &type = base.getType();
-
-            TReflection::TNameToIndex::const_iterator it = reflection.nameToIndex.find(name.c_str());
-            if (it == reflection.nameToIndex.end()) {
-                reflection.nameToIndex[name.c_str()] = (int)reflection.indexToPipeOutput.size();
-                reflection.indexToPipeOutput.push_back(TObjectReflection(name.c_str(), type, 0, mapToGlType(type), 0, 0));
+                TString baseName;
+                if (type.getBasicType() == EbtBlock) {
+                    baseName = anonymous ? TString() : type.getTypeName();
+                } else {
+                    baseName = anonymous ? TString() : name;
+                }
 
-                EShLanguageMask& stages = reflection.indexToPipeOutput.back().stages;
-                stages = static_cast<EShLanguageMask>(stages | 1 << intermediate.getStage());
+                // by convention if this is an arrayed block we ignore the array in the reflection
+                if (type.isArray()) {
+                    blowUpIOAggregate(input, baseName, TType(type, 0));
+                } else {               
+                    blowUpIOAggregate(input, baseName, type);
+                }
             } else {
-                EShLanguageMask& stages = reflection.indexToPipeOutput[it->second].stages;
-                stages = static_cast<EShLanguageMask>(stages | 1 << intermediate.getStage());
+                TReflection::TNameToIndex::const_iterator it = reflection.nameToIndex.find(name.c_str());
+                if (it == reflection.nameToIndex.end()) {
+                    reflection.nameToIndex[name.c_str()] = (int)ioItems.size();
+                    ioItems.push_back(TObjectReflection(name.c_str(), type, 0, mapToGlType(type), 0, 0));
+
+                    EShLanguageMask& stages = ioItems.back().stages;
+                    stages = static_cast<EShLanguageMask>(stages | 1 << intermediate.getStage());
+                } else {
+                    EShLanguageMask& stages = ioItems[it->second].stages;
+                    stages = static_cast<EShLanguageMask>(stages | 1 << intermediate.getStage());
+                }
             }
         }
     }
@@ -473,6 +473,67 @@ public:
             }
         }
     }
+    
+    // similar to blowUpActiveAggregate, but with simpler rules and no dereferences to follow.
+    void blowUpIOAggregate(bool input, const TString &baseName, const TType &type)
+    {
+        TString name = baseName;
+
+        // if the type is still too coarse a granularity, this is still an aggregate to expand, expand it...
+        if (! isReflectionGranularity(type)) {
+            if (type.isArray()) {
+                // Visit all the indices of this array, and for each one,
+                // fully explode the remaining aggregate to dereference
+                for (int i = 0; i < std::max(type.getOuterArraySize(), 1); ++i) {
+                    TString newBaseName = name;
+                    newBaseName.append(TString("[") + String(i) + "]");
+                    TType derefType(type, 0);
+
+                    blowUpIOAggregate(input, newBaseName, derefType);
+                }
+            } else {
+                // Visit all members of this aggregate, and for each one,
+                // fully explode the remaining aggregate to dereference
+                const TTypeList& typeList = *type.getStruct();
+
+                for (int i = 0; i < (int)typeList.size(); ++i) {
+                    TString newBaseName = name;
+                    if (newBaseName.size() > 0)
+                        newBaseName.append(".");
+                    newBaseName.append(typeList[i].type->getFieldName());
+                    TType derefType(type, i);
+
+                    blowUpIOAggregate(input, newBaseName, derefType);
+                }
+            }
+
+            // it was all completed in the recursive calls above
+            return;
+        }
+
+        if ((reflection.options & EShReflectionBasicArraySuffix) && type.isArray()) {
+            name.append(TString("[0]"));
+        }
+
+        TReflection::TMapIndexToReflection &ioItems =
+            input ? reflection.indexToPipeInput : reflection.indexToPipeOutput;
+
+        std::string namespacedName = input ? "in " : "out ";
+        namespacedName += name.c_str();
+
+        TReflection::TNameToIndex::const_iterator it = reflection.nameToIndex.find(namespacedName);
+        if (it == reflection.nameToIndex.end()) {
+            reflection.nameToIndex[namespacedName] = (int)ioItems.size();
+            ioItems.push_back(
+                TObjectReflection(name.c_str(), type, 0, mapToGlType(type), mapToGlArraySize(type), 0));
+
+            EShLanguageMask& stages = ioItems.back().stages;
+            stages = static_cast<EShLanguageMask>(stages | 1 << intermediate.getStage());
+        } else {
+            EShLanguageMask& stages = ioItems[it->second].stages;
+            stages = static_cast<EShLanguageMask>(stages | 1 << intermediate.getStage());
+        }
+    }
 
     // Add a uniform dereference where blocks/struct/arrays are involved in the access.
     // Handles the situation where the left node is at the correct or too coarse a
@@ -1027,11 +1088,9 @@ void TReflectionTraverser::visitSymbol(TIntermSymbol* base)
     if (base->getQualifier().storage == EvqUniform)
         addUniform(*base);
 
-    if (intermediate.getStage() == reflection.firstStage && base->getQualifier().isPipeInput())
-        addPipeInput(*base);
-
-    if (intermediate.getStage() == reflection.lastStage && base->getQualifier().isPipeOutput())
-        addPipeOutput(*base);
+    if ((intermediate.getStage() == reflection.firstStage && base->getQualifier().isPipeInput()) ||
+        (intermediate.getStage() == reflection.lastStage && base->getQualifier().isPipeOutput()))
+        addPipeIOVariable(*base);
 }
 
 //

3rdparty/glslang/glslang/Public/ShaderLang.h

@@ -246,6 +246,7 @@ typedef enum {
     EShReflectionIntermediateIO    = (1 << 2), // reflect inputs and outputs to program, even with no vertex shader
     EShReflectionSeparateBuffers   = (1 << 3), // buffer variables and buffer blocks are reflected separately
     EShReflectionAllBlockVariables = (1 << 4), // reflect all variables in blocks, even if they are inactive
+    EShReflectionUnwrapIOBlocks    = (1 << 5), // unwrap input/output blocks the same as with uniform blocks
 } EShReflectionOptions;
 
 //

3rdparty/glslang/gtests/AST.FromFile.cpp

@@ -96,6 +96,7 @@ INSTANTIATE_TEST_CASE_P(
         "cppNest.vert",
         "cppBad.vert",
         "cppBad2.vert",
+        "cppBad3.vert",
         "cppComplexExpr.vert",
         "cppDeepNest.frag",
         "cppPassMacroName.frag",

3rdparty/glslang/hlsl/hlslGrammar.cpp

@@ -1197,7 +1197,7 @@ bool HlslGrammar::acceptSamplerTypeDX9(TType &type)
     TSampler sampler;
     sampler.set(txType.getBasicType(), dim, false, isShadow, false);
 
-	if (!parseContext.setTextureReturnType(sampler, txType, token.loc))
+    if (!parseContext.setTextureReturnType(sampler, txType, token.loc))
         return false;
 
     type.shallowCopy(TType(sampler, EvqUniform, arraySizes));
@@ -1489,7 +1489,7 @@ bool HlslGrammar::acceptType(TType& type, TIntermNode*& nodeList)
     case EHTokSampler3d:              // ...
     case EHTokSamplerCube:            // ...
         if (parseContext.hlslDX9Compatible())
-			return acceptSamplerTypeDX9(type);
+            return acceptSamplerTypeDX9(type);
         else
             return acceptSamplerType(type);
         break;

3rdparty/glslang/hlsl/hlslParseHelper.cpp

@@ -3255,8 +3255,8 @@ void HlslParseContext::decomposeStructBufferMethods(const TSourceLoc& loc, TInte
     if (argAggregate) {
         if (argAggregate->getSequence().empty())
             return;
-		if (argAggregate->getSequence()[0])
-	        bufferObj = argAggregate->getSequence()[0]->getAsTyped();
+        if (argAggregate->getSequence()[0])
+            bufferObj = argAggregate->getSequence()[0]->getAsTyped();
     } else {
         bufferObj = arguments->getAsSymbolNode();
     }
@@ -3756,7 +3756,7 @@ void HlslParseContext::decomposeSampleMethods(const TSourceLoc& loc, TIntermType
                 return;
         } else {
             if (argAggregate->getSequence().size() == 0 || 
-				argAggregate->getSequence()[0] == nullptr ||
+                argAggregate->getSequence()[0] == nullptr ||
                 argAggregate->getSequence()[0]->getAsTyped()->getBasicType() != EbtSampler)
                 return;
         }

3rdparty/spirv-cross/reference/opt/shaders-msl/asm/tesc/tess-fixed-input-array-builtin-array.invalid.asm.tesc

@@ -0,0 +1,71 @@
+#pragma clang diagnostic ignored "-Wmissing-prototypes"
+
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+struct VertexOutput
+{
+    float4 pos;
+    float2 uv;
+};
+
+struct VertexOutput_1
+{
+    float2 uv;
+};
+
+struct HSOut
+{
+    float2 uv;
+};
+
+struct main0_out
+{
+    HSOut _entryPointOutput;
+    float4 gl_Position;
+};
+
+struct main0_in
+{
+    float2 VertexOutput_uv [[attribute(0)]];
+    float4 gl_Position [[attribute(1)]];
+};
+
+// Implementation of an array copy function to cover GLSL's ability to copy an array via assignment.
+template<typename T, uint N>
+void spvArrayCopyFromStack1(thread T (&dst)[N], thread const T (&src)[N])
+{
+    for (uint i = 0; i < N; dst[i] = src[i], i++);
+}
+
+template<typename T, uint N>
+void spvArrayCopyFromConstant1(thread T (&dst)[N], constant T (&src)[N])
+{
+    for (uint i = 0; i < N; dst[i] = src[i], i++);
+}
+
+kernel void main0(main0_in in [[stage_in]], uint gl_InvocationID [[thread_index_in_threadgroup]], uint gl_PrimitiveID [[threadgroup_position_in_grid]], device main0_out* spvOut [[buffer(28)]], device uint* spvIndirectParams [[buffer(29)]], device MTLTriangleTessellationFactorsHalf* spvTessLevel [[buffer(26)]], threadgroup main0_in* gl_in [[threadgroup(0)]])
+{
+    device main0_out* gl_out = &spvOut[gl_PrimitiveID * 3];
+    if (gl_InvocationID < spvIndirectParams[0])
+        gl_in[gl_InvocationID] = in;
+    threadgroup_barrier(mem_flags::mem_threadgroup);
+    VertexOutput _223[3] = { VertexOutput{ gl_in[0].gl_Position, gl_in[0].VertexOutput_uv }, VertexOutput{ gl_in[1].gl_Position, gl_in[1].VertexOutput_uv }, VertexOutput{ gl_in[2].gl_Position, gl_in[2].VertexOutput_uv } };
+    VertexOutput param[3];
+    spvArrayCopyFromStack1(param, _223);
+    gl_out[gl_InvocationID].gl_Position = param[gl_InvocationID].pos;
+    gl_out[gl_InvocationID]._entryPointOutput.uv = param[gl_InvocationID].uv;
+    threadgroup_barrier(mem_flags::mem_device);
+    if (int(gl_InvocationID) == 0)
+    {
+        float2 _174 = float2(1.0) + gl_in[0].VertexOutput_uv;
+        float _175 = _174.x;
+        spvTessLevel[gl_PrimitiveID].edgeTessellationFactor[0] = half(_175);
+        spvTessLevel[gl_PrimitiveID].edgeTessellationFactor[1] = half(_175);
+        spvTessLevel[gl_PrimitiveID].edgeTessellationFactor[2] = half(_175);
+        spvTessLevel[gl_PrimitiveID].insideTessellationFactor = half(_175);
+    }
+}
+

3rdparty/spirv-cross/reference/opt/shaders-msl/desktop-only/tesc/arrayed-output.desktop.sso.tesc

@@ -0,0 +1,42 @@
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+struct main0_out
+{
+    float3 vVertex;
+};
+
+struct main0_patchOut
+{
+    float3 vPatch[2];
+};
+
+struct main0_in
+{
+    float3 vInput [[attribute(0)]];
+};
+
+kernel void main0(main0_in in [[stage_in]], uint gl_InvocationID [[thread_index_in_threadgroup]], uint gl_PrimitiveID [[threadgroup_position_in_grid]], device main0_out* spvOut [[buffer(28)]], device uint* spvIndirectParams [[buffer(29)]], device main0_patchOut* spvPatchOut [[buffer(27)]], device MTLQuadTessellationFactorsHalf* spvTessLevel [[buffer(26)]], threadgroup main0_in* gl_in [[threadgroup(0)]])
+{
+    device main0_out* gl_out = &spvOut[gl_PrimitiveID * 4];
+    device main0_patchOut& patchOut = spvPatchOut[gl_PrimitiveID];
+    if (gl_InvocationID < spvIndirectParams[0])
+        gl_in[gl_InvocationID] = in;
+    threadgroup_barrier(mem_flags::mem_threadgroup);
+    gl_out[gl_InvocationID].vVertex = gl_in[gl_InvocationID].vInput + gl_in[gl_InvocationID ^ 1].vInput;
+    threadgroup_barrier(mem_flags::mem_device);
+    if (gl_InvocationID == 0)
+    {
+        patchOut.vPatch[0] = float3(10.0);
+        patchOut.vPatch[1] = float3(20.0);
+        spvTessLevel[gl_PrimitiveID].edgeTessellationFactor[0] = half(1.0);
+        spvTessLevel[gl_PrimitiveID].edgeTessellationFactor[1] = half(2.0);
+        spvTessLevel[gl_PrimitiveID].edgeTessellationFactor[2] = half(3.0);
+        spvTessLevel[gl_PrimitiveID].edgeTessellationFactor[3] = half(4.0);
+        spvTessLevel[gl_PrimitiveID].insideTessellationFactor[0] = half(1.0);
+        spvTessLevel[gl_PrimitiveID].insideTessellationFactor[1] = half(2.0);
+    }
+}
+

3rdparty/spirv-cross/reference/opt/shaders-msl/desktop-only/tesc/basic.desktop.sso.tesc

@@ -0,0 +1,37 @@
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+struct main0_out
+{
+    float4 gl_Position;
+};
+
+struct main0_patchOut
+{
+    float3 vFoo;
+};
+
+struct main0_in
+{
+    float4 gl_Position [[attribute(0)]];
+};
+
+kernel void main0(main0_in in [[stage_in]], uint gl_InvocationID [[thread_index_in_threadgroup]], uint gl_PrimitiveID [[threadgroup_position_in_grid]], device main0_out* spvOut [[buffer(28)]], device uint* spvIndirectParams [[buffer(29)]], device main0_patchOut* spvPatchOut [[buffer(27)]], device MTLQuadTessellationFactorsHalf* spvTessLevel [[buffer(26)]], threadgroup main0_in* gl_in [[threadgroup(0)]])
+{
+    device main0_out* gl_out = &spvOut[gl_PrimitiveID * 1];
+    device main0_patchOut& patchOut = spvPatchOut[gl_PrimitiveID];
+    if (gl_InvocationID < spvIndirectParams[0])
+        gl_in[gl_InvocationID] = in;
+    threadgroup_barrier(mem_flags::mem_threadgroup);
+    spvTessLevel[gl_PrimitiveID].insideTessellationFactor[0] = half(8.8999996185302734375);
+    spvTessLevel[gl_PrimitiveID].insideTessellationFactor[1] = half(6.900000095367431640625);
+    spvTessLevel[gl_PrimitiveID].edgeTessellationFactor[0] = half(8.8999996185302734375);
+    spvTessLevel[gl_PrimitiveID].edgeTessellationFactor[1] = half(6.900000095367431640625);
+    spvTessLevel[gl_PrimitiveID].edgeTessellationFactor[2] = half(3.900000095367431640625);
+    spvTessLevel[gl_PrimitiveID].edgeTessellationFactor[3] = half(4.900000095367431640625);
+    patchOut.vFoo = float3(1.0);
+    gl_out[gl_InvocationID].gl_Position = gl_in[0].gl_Position + gl_in[1].gl_Position;
+}
+

3rdparty/spirv-cross/reference/opt/shaders-msl/desktop-only/tesc/struct-copy.desktop.sso.tesc

@@ -0,0 +1,40 @@
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+struct Boo
+{
+    float3 a;
+    float3 b;
+};
+
+struct main0_out
+{
+    Boo vVertex;
+};
+
+struct main0_in
+{
+    float3 Boo_a [[attribute(0)]];
+    float3 Boo_b [[attribute(1)]];
+};
+
+kernel void main0(main0_in in [[stage_in]], uint gl_InvocationID [[thread_index_in_threadgroup]], uint gl_PrimitiveID [[threadgroup_position_in_grid]], device main0_out* spvOut [[buffer(28)]], device uint* spvIndirectParams [[buffer(29)]], device MTLQuadTessellationFactorsHalf* spvTessLevel [[buffer(26)]], threadgroup main0_in* gl_in [[threadgroup(0)]])
+{
+    device main0_out* gl_out = &spvOut[gl_PrimitiveID * 4];
+    if (gl_InvocationID < spvIndirectParams[0])
+        gl_in[gl_InvocationID] = in;
+    threadgroup_barrier(mem_flags::mem_threadgroup);
+    Boo vInput_24;
+    vInput_24.a = gl_in[gl_InvocationID].Boo_a;
+    vInput_24.b = gl_in[gl_InvocationID].Boo_b;
+    gl_out[gl_InvocationID].vVertex = vInput_24;
+    spvTessLevel[gl_PrimitiveID].edgeTessellationFactor[0] = half(1.0);
+    spvTessLevel[gl_PrimitiveID].edgeTessellationFactor[1] = half(2.0);
+    spvTessLevel[gl_PrimitiveID].edgeTessellationFactor[2] = half(3.0);
+    spvTessLevel[gl_PrimitiveID].edgeTessellationFactor[3] = half(4.0);
+    spvTessLevel[gl_PrimitiveID].insideTessellationFactor[0] = half(1.0);
+    spvTessLevel[gl_PrimitiveID].insideTessellationFactor[1] = half(2.0);
+}
+

3rdparty/spirv-cross/reference/opt/shaders-msl/tesc/basic.tesc

@@ -0,0 +1,22 @@
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+struct main0_patchOut
+{
+    float3 vFoo;
+};
+
+kernel void main0(uint gl_InvocationID [[thread_index_in_threadgroup]], uint gl_PrimitiveID [[threadgroup_position_in_grid]], device uint* spvIndirectParams [[buffer(29)]], device main0_patchOut* spvPatchOut [[buffer(27)]], device MTLQuadTessellationFactorsHalf* spvTessLevel [[buffer(26)]])
+{
+    device main0_patchOut& patchOut = spvPatchOut[gl_PrimitiveID];
+    spvTessLevel[gl_PrimitiveID].insideTessellationFactor[0] = half(8.8999996185302734375);
+    spvTessLevel[gl_PrimitiveID].insideTessellationFactor[1] = half(6.900000095367431640625);
+    spvTessLevel[gl_PrimitiveID].edgeTessellationFactor[0] = half(8.8999996185302734375);
+    spvTessLevel[gl_PrimitiveID].edgeTessellationFactor[1] = half(6.900000095367431640625);
+    spvTessLevel[gl_PrimitiveID].edgeTessellationFactor[2] = half(3.900000095367431640625);
+    spvTessLevel[gl_PrimitiveID].edgeTessellationFactor[3] = half(4.900000095367431640625);
+    patchOut.vFoo = float3(1.0);
+}
+

3rdparty/spirv-cross/reference/opt/shaders-msl/tesc/water_tess.tesc

@@ -0,0 +1,91 @@
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+struct UBO
+{
+    float4 uScale;
+    float3 uCamPos;
+    float2 uPatchSize;
+    float2 uMaxTessLevel;
+    float uDistanceMod;
+    float4 uFrustum[6];
+};
+
+struct main0_patchOut
+{
+    float2 vOutPatchPosBase;
+    float4 vPatchLods;
+};
+
+struct main0_in
+{
+    float2 vPatchPosBase [[attribute(0)]];
+};
+
+kernel void main0(main0_in in [[stage_in]], constant UBO& _41 [[buffer(0)]], uint gl_InvocationID [[thread_index_in_threadgroup]], uint gl_PrimitiveID [[threadgroup_position_in_grid]], device uint* spvIndirectParams [[buffer(29)]], device main0_patchOut* spvPatchOut [[buffer(27)]], device MTLQuadTessellationFactorsHalf* spvTessLevel [[buffer(26)]], threadgroup main0_in* gl_in [[threadgroup(0)]])
+{
+    device main0_patchOut& patchOut = spvPatchOut[gl_PrimitiveID];
+    if (gl_InvocationID < spvIndirectParams[0])
+        gl_in[gl_InvocationID] = in;
+    threadgroup_barrier(mem_flags::mem_threadgroup);
+    float2 _430 = (gl_in[0].vPatchPosBase - float2(10.0)) * _41.uScale.xy;
+    float2 _440 = ((gl_in[0].vPatchPosBase + _41.uPatchSize) + float2(10.0)) * _41.uScale.xy;
+    float3 _445 = float3(_430.x, -10.0, _430.y);
+    float3 _450 = float3(_440.x, 10.0, _440.y);
+    float4 _466 = float4((_445 + _450) * 0.5, 1.0);
+    float3 _513 = float3(length(_450 - _445) * (-0.5));
+    bool _515 = any(float3(dot(_41.uFrustum[0], _466), dot(_41.uFrustum[1], _466), dot(_41.uFrustum[2], _466)) <= _513);
+    bool _525;
+    if (!_515)
+    {
+        _525 = any(float3(dot(_41.uFrustum[3], _466), dot(_41.uFrustum[4], _466), dot(_41.uFrustum[5], _466)) <= _513);
+    }
+    else
+    {
+        _525 = _515;
+    }
+    if (!(!_525))
+    {
+        spvTessLevel[gl_PrimitiveID].edgeTessellationFactor[0] = half(-1.0);
+        spvTessLevel[gl_PrimitiveID].edgeTessellationFactor[1] = half(-1.0);
+        spvTessLevel[gl_PrimitiveID].edgeTessellationFactor[2] = half(-1.0);
+        spvTessLevel[gl_PrimitiveID].edgeTessellationFactor[3] = half(-1.0);
+        spvTessLevel[gl_PrimitiveID].insideTessellationFactor[0] = half(-1.0);
+        spvTessLevel[gl_PrimitiveID].insideTessellationFactor[1] = half(-1.0);
+    }
+    else
+    {
+        patchOut.vOutPatchPosBase = gl_in[0].vPatchPosBase;
+        float2 _678 = (gl_in[0].vPatchPosBase + (float2(-0.5) * _41.uPatchSize)) * _41.uScale.xy;
+        float2 _706 = (gl_in[0].vPatchPosBase + (float2(0.5, -0.5) * _41.uPatchSize)) * _41.uScale.xy;
+        float _725 = fast::clamp(log2((length(_41.uCamPos - float3(_706.x, 0.0, _706.y)) + 9.9999997473787516355514526367188e-05) * _41.uDistanceMod), 0.0, _41.uMaxTessLevel.x);
+        float2 _734 = (gl_in[0].vPatchPosBase + (float2(1.5, -0.5) * _41.uPatchSize)) * _41.uScale.xy;
+        float2 _762 = (gl_in[0].vPatchPosBase + (float2(-0.5, 0.5) * _41.uPatchSize)) * _41.uScale.xy;
+        float _781 = fast::clamp(log2((length(_41.uCamPos - float3(_762.x, 0.0, _762.y)) + 9.9999997473787516355514526367188e-05) * _41.uDistanceMod), 0.0, _41.uMaxTessLevel.x);
+        float2 _790 = (gl_in[0].vPatchPosBase + (float2(0.5) * _41.uPatchSize)) * _41.uScale.xy;
+        float _809 = fast::clamp(log2((length(_41.uCamPos - float3(_790.x, 0.0, _790.y)) + 9.9999997473787516355514526367188e-05) * _41.uDistanceMod), 0.0, _41.uMaxTessLevel.x);
+        float2 _818 = (gl_in[0].vPatchPosBase + (float2(1.5, 0.5) * _41.uPatchSize)) * _41.uScale.xy;
+        float _837 = fast::clamp(log2((length(_41.uCamPos - float3(_818.x, 0.0, _818.y)) + 9.9999997473787516355514526367188e-05) * _41.uDistanceMod), 0.0, _41.uMaxTessLevel.x);
+        float2 _846 = (gl_in[0].vPatchPosBase + (float2(-0.5, 1.5) * _41.uPatchSize)) * _41.uScale.xy;
+        float2 _874 = (gl_in[0].vPatchPosBase + (float2(0.5, 1.5) * _41.uPatchSize)) * _41.uScale.xy;
+        float _893 = fast::clamp(log2((length(_41.uCamPos - float3(_874.x, 0.0, _874.y)) + 9.9999997473787516355514526367188e-05) * _41.uDistanceMod), 0.0, _41.uMaxTessLevel.x);
+        float2 _902 = (gl_in[0].vPatchPosBase + (float2(1.5) * _41.uPatchSize)) * _41.uScale.xy;
+        float _612 = dot(float4(_781, _809, fast::clamp(log2((length(_41.uCamPos - float3(_846.x, 0.0, _846.y)) + 9.9999997473787516355514526367188e-05) * _41.uDistanceMod), 0.0, _41.uMaxTessLevel.x), _893), float4(0.25));
+        float _618 = dot(float4(fast::clamp(log2((length(_41.uCamPos - float3(_678.x, 0.0, _678.y)) + 9.9999997473787516355514526367188e-05) * _41.uDistanceMod), 0.0, _41.uMaxTessLevel.x), _725, _781, _809), float4(0.25));
+        float _624 = dot(float4(_725, fast::clamp(log2((length(_41.uCamPos - float3(_734.x, 0.0, _734.y)) + 9.9999997473787516355514526367188e-05) * _41.uDistanceMod), 0.0, _41.uMaxTessLevel.x), _809, _837), float4(0.25));
+        float _630 = dot(float4(_809, _837, _893, fast::clamp(log2((length(_41.uCamPos - float3(_902.x, 0.0, _902.y)) + 9.9999997473787516355514526367188e-05) * _41.uDistanceMod), 0.0, _41.uMaxTessLevel.x)), float4(0.25));
+        float4 _631 = float4(_612, _618, _624, _630);
+        patchOut.vPatchLods = _631;
+        float4 _928 = exp2(-fast::min(_631, _631.yzwx)) * _41.uMaxTessLevel.y;
+        spvTessLevel[gl_PrimitiveID].edgeTessellationFactor[0] = half(_928.x);
+        spvTessLevel[gl_PrimitiveID].edgeTessellationFactor[1] = half(_928.y);
+        spvTessLevel[gl_PrimitiveID].edgeTessellationFactor[2] = half(_928.z);
+        spvTessLevel[gl_PrimitiveID].edgeTessellationFactor[3] = half(_928.w);
+        float _935 = _41.uMaxTessLevel.y * exp2(-fast::min(fast::min(fast::min(_612, _618), fast::min(_624, _630)), _809));
+        spvTessLevel[gl_PrimitiveID].insideTessellationFactor[0] = half(_935);
+        spvTessLevel[gl_PrimitiveID].insideTessellationFactor[1] = half(_935);
+    }
+}
+

3rdparty/spirv-cross/reference/opt/shaders/asm/comp/recompile-block-naming.asm.comp

@@ -0,0 +1,35 @@
+#version 450
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+
+layout(binding = 0, std430) buffer MyFirstBuffer
+{
+    uint _data[];
+} MyFirstBuffer_1;
+
+layout(binding = 0, std430) buffer MySecondBuffer
+{
+    uint _data[];
+} MySecondBuffer_1;
+
+layout(binding = 0, std430) buffer MyThirdBuffer
+{
+    uint _data[];
+} MyThirdBuffer_1;
+
+void main()
+{
+    uint _105 = MyFirstBuffer_1._data[0];
+    uint _109 = MyFirstBuffer_1._data[1];
+    uint _113 = MyFirstBuffer_1._data[2];
+    uint _117 = MyFirstBuffer_1._data[3];
+    uint _122 = MySecondBuffer_1._data[1];
+    uint _126 = MySecondBuffer_1._data[2];
+    uint _130 = MySecondBuffer_1._data[3];
+    uint _134 = MySecondBuffer_1._data[4];
+    uvec4 _140 = uvec4(_105, _109, _113, _117) + uvec4(_122, _126, _130, _134);
+    MyThirdBuffer_1._data[0] = _140.x;
+    MyThirdBuffer_1._data[1] = _140.y;
+    MyThirdBuffer_1._data[2] = _140.z;
+    MyThirdBuffer_1._data[3] = _140.w;
+}
+

3rdparty/spirv-cross/reference/shaders-msl/asm/tesc/tess-fixed-input-array-builtin-array.invalid.asm.tesc

@@ -0,0 +1,111 @@
+#pragma clang diagnostic ignored "-Wmissing-prototypes"
+
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+struct VertexOutput
+{
+    float4 pos;
+    float2 uv;
+};
+
+struct HSOut
+{
+    float4 pos;
+    float2 uv;
+};
+
+struct HSConstantOut
+{
+    float EdgeTess[3];
+    float InsideTess;
+};
+
+struct VertexOutput_1
+{
+    float2 uv;
+};
+
+struct HSOut_1
+{
+    float2 uv;
+};
+
+struct main0_out
+{
+    HSOut_1 _entryPointOutput;
+    float4 gl_Position;
+};
+
+struct main0_in
+{
+    float2 VertexOutput_uv [[attribute(0)]];
+    float4 gl_Position [[attribute(1)]];
+};
+
+// Implementation of an array copy function to cover GLSL's ability to copy an array via assignment.
+template<typename T, uint N>
+void spvArrayCopyFromStack1(thread T (&dst)[N], thread const T (&src)[N])
+{
+    for (uint i = 0; i < N; dst[i] = src[i], i++);
+}
+
+template<typename T, uint N>
+void spvArrayCopyFromConstant1(thread T (&dst)[N], constant T (&src)[N])
+{
+    for (uint i = 0; i < N; dst[i] = src[i], i++);
+}
+
+HSOut _hs_main(thread const VertexOutput (&p)[3], thread const uint& i)
+{
+    HSOut _output;
+    _output.pos = p[i].pos;
+    _output.uv = p[i].uv;
+    return _output;
+}
+
+HSConstantOut PatchHS(thread const VertexOutput (&_patch)[3])
+{
+    HSConstantOut _output;
+    _output.EdgeTess[0] = (float2(1.0) + _patch[0].uv).x;
+    _output.EdgeTess[1] = (float2(1.0) + _patch[0].uv).x;
+    _output.EdgeTess[2] = (float2(1.0) + _patch[0].uv).x;
+    _output.InsideTess = (float2(1.0) + _patch[0].uv).x;
+    return _output;
+}
+
+kernel void main0(main0_in in [[stage_in]], uint gl_InvocationID [[thread_index_in_threadgroup]], uint gl_PrimitiveID [[threadgroup_position_in_grid]], device main0_out* spvOut [[buffer(28)]], device uint* spvIndirectParams [[buffer(29)]], device MTLTriangleTessellationFactorsHalf* spvTessLevel [[buffer(26)]], threadgroup main0_in* gl_in [[threadgroup(0)]])
+{
+    device main0_out* gl_out = &spvOut[gl_PrimitiveID * 3];
+    if (gl_InvocationID < spvIndirectParams[0])
+        gl_in[gl_InvocationID] = in;
+    threadgroup_barrier(mem_flags::mem_threadgroup);
+    VertexOutput p[3];
+    p[0].pos = gl_in[0].gl_Position;
+    p[0].uv = gl_in[0].VertexOutput_uv;
+    p[1].pos = gl_in[1].gl_Position;
+    p[1].uv = gl_in[1].VertexOutput_uv;
+    p[2].pos = gl_in[2].gl_Position;
+    p[2].uv = gl_in[2].VertexOutput_uv;
+    uint i = gl_InvocationID;
+    VertexOutput param[3];
+    spvArrayCopyFromStack1(param, p);
+    uint param_1 = i;
+    HSOut flattenTemp = _hs_main(param, param_1);
+    gl_out[gl_InvocationID].gl_Position = flattenTemp.pos;
+    gl_out[gl_InvocationID]._entryPointOutput.uv = flattenTemp.uv;
+    threadgroup_barrier(mem_flags::mem_device);
+    if (int(gl_InvocationID) == 0)
+    {
+        VertexOutput param_2[3];
+        spvArrayCopyFromStack1(param_2, p);
+        HSConstantOut _patchConstantResult = PatchHS(param_2);
+        spvTessLevel[gl_PrimitiveID].edgeTessellationFactor[0] = half(_patchConstantResult.EdgeTess[0]);
+        spvTessLevel[gl_PrimitiveID].edgeTessellationFactor[1] = half(_patchConstantResult.EdgeTess[1]);
+        spvTessLevel[gl_PrimitiveID].edgeTessellationFactor[2] = half(_patchConstantResult.EdgeTess[2]);
+        spvTessLevel[gl_PrimitiveID].insideTessellationFactor = half(_patchConstantResult.InsideTess);
+    }
+}
+

3rdparty/spirv-cross/reference/shaders-msl/desktop-only/tesc/arrayed-output.desktop.sso.tesc

@@ -0,0 +1,42 @@
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+struct main0_out
+{
+    float3 vVertex;
+};
+
+struct main0_patchOut
+{
+    float3 vPatch[2];
+};
+
+struct main0_in
+{
+    float3 vInput [[attribute(0)]];
+};
+
+kernel void main0(main0_in in [[stage_in]], uint gl_InvocationID [[thread_index_in_threadgroup]], uint gl_PrimitiveID [[threadgroup_position_in_grid]], device main0_out* spvOut [[buffer(28)]], device uint* spvIndirectParams [[buffer(29)]], device main0_patchOut* spvPatchOut [[buffer(27)]], device MTLQuadTessellationFactorsHalf* spvTessLevel [[buffer(26)]], threadgroup main0_in* gl_in [[threadgroup(0)]])
+{
+    device main0_out* gl_out = &spvOut[gl_PrimitiveID * 4];
+    device main0_patchOut& patchOut = spvPatchOut[gl_PrimitiveID];
+    if (gl_InvocationID < spvIndirectParams[0])
+        gl_in[gl_InvocationID] = in;
+    threadgroup_barrier(mem_flags::mem_threadgroup);
+    gl_out[gl_InvocationID].vVertex = gl_in[gl_InvocationID].vInput + gl_in[gl_InvocationID ^ 1].vInput;
+    threadgroup_barrier(mem_flags::mem_device);
+    if (gl_InvocationID == 0)
+    {
+        patchOut.vPatch[0] = float3(10.0);
+        patchOut.vPatch[1] = float3(20.0);
+        spvTessLevel[gl_PrimitiveID].edgeTessellationFactor[0] = half(1.0);
+        spvTessLevel[gl_PrimitiveID].edgeTessellationFactor[1] = half(2.0);
+        spvTessLevel[gl_PrimitiveID].edgeTessellationFactor[2] = half(3.0);
+        spvTessLevel[gl_PrimitiveID].edgeTessellationFactor[3] = half(4.0);
+        spvTessLevel[gl_PrimitiveID].insideTessellationFactor[0] = half(1.0);
+        spvTessLevel[gl_PrimitiveID].insideTessellationFactor[1] = half(2.0);
+    }
+}
+

3rdparty/spirv-cross/reference/shaders-msl/desktop-only/tesc/basic.desktop.sso.tesc

@@ -0,0 +1,44 @@
+#pragma clang diagnostic ignored "-Wmissing-prototypes"
+
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+struct main0_out
+{
+    float4 gl_Position;
+};
+
+struct main0_patchOut
+{
+    float3 vFoo;
+};
+
+struct main0_in
+{
+    float4 gl_Position [[attribute(0)]];
+};
+
+void set_position(device main0_out* thread & gl_out, thread uint& gl_InvocationID, threadgroup main0_in* thread & gl_in)
+{
+    gl_out[gl_InvocationID].gl_Position = gl_in[0].gl_Position + gl_in[1].gl_Position;
+}
+
+kernel void main0(main0_in in [[stage_in]], uint gl_InvocationID [[thread_index_in_threadgroup]], uint gl_PrimitiveID [[threadgroup_position_in_grid]], device main0_out* spvOut [[buffer(28)]], device uint* spvIndirectParams [[buffer(29)]], device main0_patchOut* spvPatchOut [[buffer(27)]], device MTLQuadTessellationFactorsHalf* spvTessLevel [[buffer(26)]], threadgroup main0_in* gl_in [[threadgroup(0)]])
+{
+    device main0_out* gl_out = &spvOut[gl_PrimitiveID * 1];
+    device main0_patchOut& patchOut = spvPatchOut[gl_PrimitiveID];
+    if (gl_InvocationID < spvIndirectParams[0])
+        gl_in[gl_InvocationID] = in;
+    threadgroup_barrier(mem_flags::mem_threadgroup);
+    spvTessLevel[gl_PrimitiveID].insideTessellationFactor[0] = half(8.8999996185302734375);
+    spvTessLevel[gl_PrimitiveID].insideTessellationFactor[1] = half(6.900000095367431640625);
+    spvTessLevel[gl_PrimitiveID].edgeTessellationFactor[0] = half(8.8999996185302734375);
+    spvTessLevel[gl_PrimitiveID].edgeTessellationFactor[1] = half(6.900000095367431640625);
+    spvTessLevel[gl_PrimitiveID].edgeTessellationFactor[2] = half(3.900000095367431640625);
+    spvTessLevel[gl_PrimitiveID].edgeTessellationFactor[3] = half(4.900000095367431640625);
+    patchOut.vFoo = float3(1.0);
+    set_position(gl_out, gl_InvocationID, gl_in);
+}
+

3rdparty/spirv-cross/reference/shaders-msl/desktop-only/tesc/struct-copy.desktop.sso.tesc

@@ -0,0 +1,40 @@
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+struct Boo
+{
+    float3 a;
+    float3 b;
+};
+
+struct main0_out
+{
+    Boo vVertex;
+};
+
+struct main0_in
+{
+    float3 Boo_a [[attribute(0)]];
+    float3 Boo_b [[attribute(1)]];
+};
+
+kernel void main0(main0_in in [[stage_in]], uint gl_InvocationID [[thread_index_in_threadgroup]], uint gl_PrimitiveID [[threadgroup_position_in_grid]], device main0_out* spvOut [[buffer(28)]], device uint* spvIndirectParams [[buffer(29)]], device MTLQuadTessellationFactorsHalf* spvTessLevel [[buffer(26)]], threadgroup main0_in* gl_in [[threadgroup(0)]])
+{
+    device main0_out* gl_out = &spvOut[gl_PrimitiveID * 4];
+    if (gl_InvocationID < spvIndirectParams[0])
+        gl_in[gl_InvocationID] = in;
+    threadgroup_barrier(mem_flags::mem_threadgroup);
+    Boo vInput_24;
+    vInput_24.a = gl_in[gl_InvocationID].Boo_a;
+    vInput_24.b = gl_in[gl_InvocationID].Boo_b;
+    gl_out[gl_InvocationID].vVertex = vInput_24;
+    spvTessLevel[gl_PrimitiveID].edgeTessellationFactor[0] = half(1.0);
+    spvTessLevel[gl_PrimitiveID].edgeTessellationFactor[1] = half(2.0);
+    spvTessLevel[gl_PrimitiveID].edgeTessellationFactor[2] = half(3.0);
+    spvTessLevel[gl_PrimitiveID].edgeTessellationFactor[3] = half(4.0);
+    spvTessLevel[gl_PrimitiveID].insideTessellationFactor[0] = half(1.0);
+    spvTessLevel[gl_PrimitiveID].insideTessellationFactor[1] = half(2.0);
+}
+

3rdparty/spirv-cross/reference/shaders-msl/tesc/basic.tesc

@@ -0,0 +1,22 @@
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+struct main0_patchOut
+{
+    float3 vFoo;
+};
+
+kernel void main0(uint gl_InvocationID [[thread_index_in_threadgroup]], uint gl_PrimitiveID [[threadgroup_position_in_grid]], device uint* spvIndirectParams [[buffer(29)]], device main0_patchOut* spvPatchOut [[buffer(27)]], device MTLQuadTessellationFactorsHalf* spvTessLevel [[buffer(26)]])
+{
+    device main0_patchOut& patchOut = spvPatchOut[gl_PrimitiveID];
+    spvTessLevel[gl_PrimitiveID].insideTessellationFactor[0] = half(8.8999996185302734375);
+    spvTessLevel[gl_PrimitiveID].insideTessellationFactor[1] = half(6.900000095367431640625);
+    spvTessLevel[gl_PrimitiveID].edgeTessellationFactor[0] = half(8.8999996185302734375);
+    spvTessLevel[gl_PrimitiveID].edgeTessellationFactor[1] = half(6.900000095367431640625);
+    spvTessLevel[gl_PrimitiveID].edgeTessellationFactor[2] = half(3.900000095367431640625);
+    spvTessLevel[gl_PrimitiveID].edgeTessellationFactor[3] = half(4.900000095367431640625);
+    patchOut.vFoo = float3(1.0);
+}
+

3rdparty/spirv-cross/reference/shaders-msl/tesc/water_tess.tesc

@@ -0,0 +1,132 @@
+#pragma clang diagnostic ignored "-Wmissing-prototypes"
+
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+struct UBO
+{
+    float4 uScale;
+    float3 uCamPos;
+    float2 uPatchSize;
+    float2 uMaxTessLevel;
+    float uDistanceMod;
+    float4 uFrustum[6];
+};
+
+struct main0_patchOut
+{
+    float2 vOutPatchPosBase;
+    float4 vPatchLods;
+};
+
+struct main0_in
+{
+    float2 vPatchPosBase [[attribute(0)]];
+};
+
+bool frustum_cull(thread const float2& p0, constant UBO& v_41)
+{
+    float2 min_xz = (p0 - float2(10.0)) * v_41.uScale.xy;
+    float2 max_xz = ((p0 + v_41.uPatchSize) + float2(10.0)) * v_41.uScale.xy;
+    float3 bb_min = float3(min_xz.x, -10.0, min_xz.y);
+    float3 bb_max = float3(max_xz.x, 10.0, max_xz.y);
+    float3 center = (bb_min + bb_max) * 0.5;
+    float radius = 0.5 * length(bb_max - bb_min);
+    float3 f0 = float3(dot(v_41.uFrustum[0], float4(center, 1.0)), dot(v_41.uFrustum[1], float4(center, 1.0)), dot(v_41.uFrustum[2], float4(center, 1.0)));
+    float3 f1 = float3(dot(v_41.uFrustum[3], float4(center, 1.0)), dot(v_41.uFrustum[4], float4(center, 1.0)), dot(v_41.uFrustum[5], float4(center, 1.0)));
+    float3 _199 = f0;
+    float _200 = radius;
+    bool _205 = any(_199 <= float3(-_200));
+    bool _215;
+    if (!_205)
+    {
+        _215 = any(f1 <= float3(-radius));
+    }
+    else
+    {
+        _215 = _205;
+    }
+    return !_215;
+}
+
+float lod_factor(thread const float2& pos_, constant UBO& v_41)
+{
+    float2 pos = pos_ * v_41.uScale.xy;
+    float3 dist_to_cam = v_41.uCamPos - float3(pos.x, 0.0, pos.y);
+    float level = log2((length(dist_to_cam) + 9.9999997473787516355514526367188e-05) * v_41.uDistanceMod);
+    return fast::clamp(level, 0.0, v_41.uMaxTessLevel.x);
+}
+
+float4 tess_level(thread const float4& lod, constant UBO& v_41)
+{
+    return exp2(-lod) * v_41.uMaxTessLevel.y;
+}
+
+float tess_level(thread const float& lod, constant UBO& v_41)
+{
+    return v_41.uMaxTessLevel.y * exp2(-lod);
+}
+
+void compute_tess_levels(thread const float2& p0, constant UBO& v_41, device float2& vOutPatchPosBase, device float4& vPatchLods, device half (&gl_TessLevelOuter)[4], device half (&gl_TessLevelInner)[2])
+{
+    vOutPatchPosBase = p0;
+    float2 param = p0 + (float2(-0.5) * v_41.uPatchSize);
+    float l00 = lod_factor(param, v_41);
+    float2 param_1 = p0 + (float2(0.5, -0.5) * v_41.uPatchSize);
+    float l10 = lod_factor(param_1, v_41);
+    float2 param_2 = p0 + (float2(1.5, -0.5) * v_41.uPatchSize);
+    float l20 = lod_factor(param_2, v_41);
+    float2 param_3 = p0 + (float2(-0.5, 0.5) * v_41.uPatchSize);
+    float l01 = lod_factor(param_3, v_41);
+    float2 param_4 = p0 + (float2(0.5) * v_41.uPatchSize);
+    float l11 = lod_factor(param_4, v_41);
+    float2 param_5 = p0 + (float2(1.5, 0.5) * v_41.uPatchSize);
+    float l21 = lod_factor(param_5, v_41);
+    float2 param_6 = p0 + (float2(-0.5, 1.5) * v_41.uPatchSize);
+    float l02 = lod_factor(param_6, v_41);
+    float2 param_7 = p0 + (float2(0.5, 1.5) * v_41.uPatchSize);
+    float l12 = lod_factor(param_7, v_41);
+    float2 param_8 = p0 + (float2(1.5) * v_41.uPatchSize);
+    float l22 = lod_factor(param_8, v_41);
+    float4 lods = float4(dot(float4(l01, l11, l02, l12), float4(0.25)), dot(float4(l00, l10, l01, l11), float4(0.25)), dot(float4(l10, l20, l11, l21), float4(0.25)), dot(float4(l11, l21, l12, l22), float4(0.25)));
+    vPatchLods = lods;
+    float4 outer_lods = fast::min(lods, lods.yzwx);
+    float4 param_9 = outer_lods;
+    float4 levels = tess_level(param_9, v_41);
+    gl_TessLevelOuter[0] = half(levels.x);
+    gl_TessLevelOuter[1] = half(levels.y);
+    gl_TessLevelOuter[2] = half(levels.z);
+    gl_TessLevelOuter[3] = half(levels.w);
+    float min_lod = fast::min(fast::min(lods.x, lods.y), fast::min(lods.z, lods.w));
+    float param_10 = fast::min(min_lod, l11);
+    float inner = tess_level(param_10, v_41);
+    gl_TessLevelInner[0] = half(inner);
+    gl_TessLevelInner[1] = half(inner);
+}
+
+kernel void main0(main0_in in [[stage_in]], constant UBO& v_41 [[buffer(0)]], uint gl_InvocationID [[thread_index_in_threadgroup]], uint gl_PrimitiveID [[threadgroup_position_in_grid]], device uint* spvIndirectParams [[buffer(29)]], device main0_patchOut* spvPatchOut [[buffer(27)]], device MTLQuadTessellationFactorsHalf* spvTessLevel [[buffer(26)]], threadgroup main0_in* gl_in [[threadgroup(0)]])
+{
+    device main0_patchOut& patchOut = spvPatchOut[gl_PrimitiveID];
+    if (gl_InvocationID < spvIndirectParams[0])
+        gl_in[gl_InvocationID] = in;
+    threadgroup_barrier(mem_flags::mem_threadgroup);
+    float2 p0 = gl_in[0].vPatchPosBase;
+    float2 param = p0;
+    if (!frustum_cull(param, v_41))
+    {
+        spvTessLevel[gl_PrimitiveID].edgeTessellationFactor[0] = half(-1.0);
+        spvTessLevel[gl_PrimitiveID].edgeTessellationFactor[1] = half(-1.0);
+        spvTessLevel[gl_PrimitiveID].edgeTessellationFactor[2] = half(-1.0);
+        spvTessLevel[gl_PrimitiveID].edgeTessellationFactor[3] = half(-1.0);
+        spvTessLevel[gl_PrimitiveID].insideTessellationFactor[0] = half(-1.0);
+        spvTessLevel[gl_PrimitiveID].insideTessellationFactor[1] = half(-1.0);
+    }
+    else
+    {
+        float2 param_1 = p0;
+        compute_tess_levels(param_1, v_41, patchOut.vOutPatchPosBase, patchOut.vPatchLods, spvTessLevel[gl_PrimitiveID].edgeTessellationFactor, spvTessLevel[gl_PrimitiveID].insideTessellationFactor);
+    }
+}
+

3rdparty/spirv-cross/reference/shaders/asm/comp/recompile-block-naming.asm.comp

@@ -0,0 +1,36 @@
+#version 450
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+
+layout(binding = 0, std430) buffer MyFirstBuffer
+{
+    uint _data[];
+} MyFirstBuffer_1;
+
+layout(binding = 0, std430) buffer MySecondBuffer
+{
+    uint _data[];
+} MySecondBuffer_1;
+
+layout(binding = 0, std430) buffer MyThirdBuffer
+{
+    uint _data[];
+} MyThirdBuffer_1;
+
+void _main()
+{
+    int byteAddrTemp = 0 >> 2;
+    uvec4 a = uvec4(MyFirstBuffer_1._data[byteAddrTemp], MyFirstBuffer_1._data[byteAddrTemp + 1], MyFirstBuffer_1._data[byteAddrTemp + 2], MyFirstBuffer_1._data[byteAddrTemp + 3]);
+    int byteAddrTemp_1 = 4 >> 2;
+    uvec4 b = uvec4(MySecondBuffer_1._data[byteAddrTemp_1], MySecondBuffer_1._data[byteAddrTemp_1 + 1], MySecondBuffer_1._data[byteAddrTemp_1 + 2], MySecondBuffer_1._data[byteAddrTemp_1 + 3]);
+    int byteAddrTemp_2 = 0 >> 2;
+    MyThirdBuffer_1._data[byteAddrTemp_2] = (a + b).x;
+    MyThirdBuffer_1._data[byteAddrTemp_2 + 1] = (a + b).y;
+    MyThirdBuffer_1._data[byteAddrTemp_2 + 2] = (a + b).z;
+    MyThirdBuffer_1._data[byteAddrTemp_2 + 3] = (a + b).w;
+}
+
+void main()
+{
+    _main();
+}
+

3rdparty/spirv-cross/shaders-msl/asm/tesc/tess-fixed-input-array-builtin-array.invalid.asm.tesc

@@ -0,0 +1,248 @@
+; SPIR-V
+; Version: 1.0
+; Generator: Khronos Glslang Reference Front End; 2
+; Bound: 162
+; Schema: 0
+               OpCapability Tessellation
+          %1 = OpExtInstImport "GLSL.std.450"
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint TessellationControl %hs_main "main" %p_pos %p_1 %i_1 %_entryPointOutput_pos %_entryPointOutput %_patchConstantOutput_EdgeTess %_patchConstantOutput_InsideTess
+               OpExecutionMode %hs_main OutputVertices 3
+               OpExecutionMode %hs_main Triangles
+               OpExecutionMode %hs_main SpacingFractionalOdd
+               OpExecutionMode %hs_main VertexOrderCw
+               OpSource HLSL 500
+               OpName %hs_main "hs_main"
+               OpName %VertexOutput "VertexOutput"
+               OpMemberName %VertexOutput 0 "pos"
+               OpMemberName %VertexOutput 1 "uv"
+               OpName %HSOut "HSOut"
+               OpMemberName %HSOut 0 "pos"
+               OpMemberName %HSOut 1 "uv"
+               OpName %_hs_main_struct_VertexOutput_vf4_vf21_3__u1_ "@hs_main(struct-VertexOutput-vf4-vf21[3];u1;"
+               OpName %p "p"
+               OpName %i "i"
+               OpName %HSConstantOut "HSConstantOut"
+               OpMemberName %HSConstantOut 0 "EdgeTess"
+               OpMemberName %HSConstantOut 1 "InsideTess"
+               OpName %PatchHS_struct_VertexOutput_vf4_vf21_3__ "PatchHS(struct-VertexOutput-vf4-vf21[3];"
+               OpName %patch "patch"
+               OpName %output "output"
+               OpName %p_0 "p"
+               OpName %p_pos "p.pos"
+               OpName %VertexOutput_0 "VertexOutput"
+               OpMemberName %VertexOutput_0 0 "uv"
+               OpName %p_1 "p"
+               OpName %i_0 "i"
+               OpName %i_1 "i"
+               OpName %flattenTemp "flattenTemp"
+               OpName %param "param"
+               OpName %param_0 "param"
+               OpName %_entryPointOutput_pos "@entryPointOutput.pos"
+               OpName %HSOut_0 "HSOut"
+               OpMemberName %HSOut_0 0 "uv"
+               OpName %_entryPointOutput "@entryPointOutput"
+               OpName %_patchConstantResult "@patchConstantResult"
+               OpName %param_1 "param"
+               OpName %_patchConstantOutput_EdgeTess "@patchConstantOutput.EdgeTess"
+               OpName %_patchConstantOutput_InsideTess "@patchConstantOutput.InsideTess"
+               OpName %output_0 "output"
+               OpDecorate %p_pos BuiltIn Position
+               OpDecorate %p_1 Location 0
+               OpDecorate %i_1 BuiltIn InvocationId
+               OpDecorate %_entryPointOutput_pos BuiltIn Position
+               OpDecorate %_entryPointOutput Location 0
+               OpDecorate %_patchConstantOutput_EdgeTess Patch
+               OpDecorate %_patchConstantOutput_EdgeTess BuiltIn TessLevelOuter
+               OpDecorate %_patchConstantOutput_InsideTess Patch
+               OpDecorate %_patchConstantOutput_InsideTess BuiltIn TessLevelInner
+       %void = OpTypeVoid
+          %3 = OpTypeFunction %void
+      %float = OpTypeFloat 32
+    %v4float = OpTypeVector %float 4
+    %v2float = OpTypeVector %float 2
+%VertexOutput = OpTypeStruct %v4float %v2float
+       %uint = OpTypeInt 32 0
+     %uint_3 = OpConstant %uint 3
+%_arr_VertexOutput_uint_3 = OpTypeArray %VertexOutput %uint_3
+%_ptr_Function__arr_VertexOutput_uint_3 = OpTypePointer Function %_arr_VertexOutput_uint_3
+%_ptr_Function_uint = OpTypePointer Function %uint
+      %HSOut = OpTypeStruct %v4float %v2float
+         %16 = OpTypeFunction %HSOut %_ptr_Function__arr_VertexOutput_uint_3 %_ptr_Function_uint
+%_arr_float_uint_3 = OpTypeArray %float %uint_3
+%HSConstantOut = OpTypeStruct %_arr_float_uint_3 %float
+         %23 = OpTypeFunction %HSConstantOut %_ptr_Function__arr_VertexOutput_uint_3
+%_ptr_Function_HSOut = OpTypePointer Function %HSOut
+        %int = OpTypeInt 32 1
+      %int_0 = OpConstant %int 0
+%_ptr_Function_v4float = OpTypePointer Function %v4float
+      %int_1 = OpConstant %int 1
+%_ptr_Function_v2float = OpTypePointer Function %v2float
+%_arr_v4float_uint_3 = OpTypeArray %v4float %uint_3
+%_ptr_Input__arr_v4float_uint_3 = OpTypePointer Input %_arr_v4float_uint_3
+      %p_pos = OpVariable %_ptr_Input__arr_v4float_uint_3 Input
+%_ptr_Input_v4float = OpTypePointer Input %v4float
+%VertexOutput_0 = OpTypeStruct %v2float
+%_arr_VertexOutput_0_uint_3 = OpTypeArray %VertexOutput_0 %uint_3
+%_ptr_Input__arr_VertexOutput_0_uint_3 = OpTypePointer Input %_arr_VertexOutput_0_uint_3
+        %p_1 = OpVariable %_ptr_Input__arr_VertexOutput_0_uint_3 Input
+%_ptr_Input_v2float = OpTypePointer Input %v2float
+      %int_2 = OpConstant %int 2
+%_ptr_Input_uint = OpTypePointer Input %uint
+        %i_1 = OpVariable %_ptr_Input_uint Input
+%_ptr_Output__arr_v4float_uint_3 = OpTypePointer Output %_arr_v4float_uint_3
+%_entryPointOutput_pos = OpVariable %_ptr_Output__arr_v4float_uint_3 Output
+%_ptr_Output_v4float = OpTypePointer Output %v4float
+    %HSOut_0 = OpTypeStruct %v2float
+%_arr_HSOut_0_uint_3 = OpTypeArray %HSOut_0 %uint_3
+%_ptr_Output__arr_HSOut_0_uint_3 = OpTypePointer Output %_arr_HSOut_0_uint_3
+%_entryPointOutput = OpVariable %_ptr_Output__arr_HSOut_0_uint_3 Output
+%_ptr_Output_v2float = OpTypePointer Output %v2float
+     %uint_2 = OpConstant %uint 2
+     %uint_1 = OpConstant %uint 1
+     %uint_0 = OpConstant %uint 0
+       %bool = OpTypeBool
+%_ptr_Function_HSConstantOut = OpTypePointer Function %HSConstantOut
+     %uint_4 = OpConstant %uint 4
+%_arr_float_uint_4 = OpTypeArray %float %uint_4
+%_ptr_Output__arr_float_uint_4 = OpTypePointer Output %_arr_float_uint_4
+%_patchConstantOutput_EdgeTess = OpVariable %_ptr_Output__arr_float_uint_4 Output
+%_ptr_Function_float = OpTypePointer Function %float
+%_ptr_Output_float = OpTypePointer Output %float
+%_arr_float_uint_2 = OpTypeArray %float %uint_2
+%_ptr_Output__arr_float_uint_2 = OpTypePointer Output %_arr_float_uint_2
+%_patchConstantOutput_InsideTess = OpVariable %_ptr_Output__arr_float_uint_2 Output
+    %float_1 = OpConstant %float 1
+    %hs_main = OpFunction %void None %3
+          %5 = OpLabel
+        %p_0 = OpVariable %_ptr_Function__arr_VertexOutput_uint_3 Function
+        %i_0 = OpVariable %_ptr_Function_uint Function
+%flattenTemp = OpVariable %_ptr_Function_HSOut Function
+      %param = OpVariable %_ptr_Function__arr_VertexOutput_uint_3 Function
+    %param_0 = OpVariable %_ptr_Function_uint Function
+%_patchConstantResult = OpVariable %_ptr_Function_HSConstantOut Function
+    %param_1 = OpVariable %_ptr_Function__arr_VertexOutput_uint_3 Function
+         %50 = OpAccessChain %_ptr_Input_v4float %p_pos %int_0
+         %51 = OpLoad %v4float %50
+         %52 = OpAccessChain %_ptr_Function_v4float %p_0 %int_0 %int_0
+               OpStore %52 %51
+         %58 = OpAccessChain %_ptr_Input_v2float %p_1 %int_0 %int_0
+         %59 = OpLoad %v2float %58
+         %60 = OpAccessChain %_ptr_Function_v2float %p_0 %int_0 %int_1
+               OpStore %60 %59
+         %61 = OpAccessChain %_ptr_Input_v4float %p_pos %int_1
+         %62 = OpLoad %v4float %61
+         %63 = OpAccessChain %_ptr_Function_v4float %p_0 %int_1 %int_0
+               OpStore %63 %62
+         %64 = OpAccessChain %_ptr_Input_v2float %p_1 %int_1 %int_0
+         %65 = OpLoad %v2float %64
+         %66 = OpAccessChain %_ptr_Function_v2float %p_0 %int_1 %int_1
+               OpStore %66 %65
+         %68 = OpAccessChain %_ptr_Input_v4float %p_pos %int_2
+         %69 = OpLoad %v4float %68
+         %70 = OpAccessChain %_ptr_Function_v4float %p_0 %int_2 %int_0
+               OpStore %70 %69
+         %71 = OpAccessChain %_ptr_Input_v2float %p_1 %int_2 %int_0
+         %72 = OpLoad %v2float %71
+         %73 = OpAccessChain %_ptr_Function_v2float %p_0 %int_2 %int_1
+               OpStore %73 %72
+         %77 = OpLoad %uint %i_1
+               OpStore %i_0 %77
+         %80 = OpLoad %_arr_VertexOutput_uint_3 %p_0
+               OpStore %param %80
+         %82 = OpLoad %uint %i_0
+               OpStore %param_0 %82
+         %83 = OpFunctionCall %HSOut %_hs_main_struct_VertexOutput_vf4_vf21_3__u1_ %param %param_0
+               OpStore %flattenTemp %83
+         %86 = OpAccessChain %_ptr_Function_v4float %flattenTemp %int_0
+         %87 = OpLoad %v4float %86
+         %94 = OpLoad %uint %i_1
+         %89 = OpAccessChain %_ptr_Output_v4float %_entryPointOutput_pos %94
+               OpStore %89 %87
+         %95 = OpAccessChain %_ptr_Function_v2float %flattenTemp %int_1
+         %96 = OpLoad %v2float %95
+         %98 = OpAccessChain %_ptr_Output_v2float %_entryPointOutput %94 %int_0
+               OpStore %98 %96
+               OpControlBarrier %uint_2 %uint_1 %uint_0
+        %102 = OpLoad %uint %i_1
+        %104 = OpIEqual %bool %102 %int_0
+               OpSelectionMerge %106 None
+               OpBranchConditional %104 %105 %106
+        %105 = OpLabel
+        %110 = OpLoad %_arr_VertexOutput_uint_3 %p_0
+               OpStore %param_1 %110
+        %111 = OpFunctionCall %HSConstantOut %PatchHS_struct_VertexOutput_vf4_vf21_3__ %param_1
+               OpStore %_patchConstantResult %111
+        %117 = OpAccessChain %_ptr_Function_float %_patchConstantResult %int_0 %int_0
+        %118 = OpLoad %float %117
+        %120 = OpAccessChain %_ptr_Output_float %_patchConstantOutput_EdgeTess %int_0
+               OpStore %120 %118
+        %121 = OpAccessChain %_ptr_Function_float %_patchConstantResult %int_0 %int_1
+        %122 = OpLoad %float %121
+        %123 = OpAccessChain %_ptr_Output_float %_patchConstantOutput_EdgeTess %int_1
+               OpStore %123 %122
+        %124 = OpAccessChain %_ptr_Function_float %_patchConstantResult %int_0 %int_2
+        %125 = OpLoad %float %124
+        %126 = OpAccessChain %_ptr_Output_float %_patchConstantOutput_EdgeTess %int_2
+               OpStore %126 %125
+        %130 = OpAccessChain %_ptr_Function_float %_patchConstantResult %int_1
+        %131 = OpLoad %float %130
+        %132 = OpAccessChain %_ptr_Output_float %_patchConstantOutput_InsideTess %int_0
+               OpStore %132 %131
+               OpBranch %106
+        %106 = OpLabel
+               OpReturn
+               OpFunctionEnd
+%_hs_main_struct_VertexOutput_vf4_vf21_3__u1_ = OpFunction %HSOut None %16
+          %p = OpFunctionParameter %_ptr_Function__arr_VertexOutput_uint_3
+          %i = OpFunctionParameter %_ptr_Function_uint
+         %20 = OpLabel
+     %output = OpVariable %_ptr_Function_HSOut Function
+         %31 = OpLoad %uint %i
+         %33 = OpAccessChain %_ptr_Function_v4float %p %31 %int_0
+         %34 = OpLoad %v4float %33
+         %35 = OpAccessChain %_ptr_Function_v4float %output %int_0
+               OpStore %35 %34
+         %37 = OpLoad %uint %i
+         %39 = OpAccessChain %_ptr_Function_v2float %p %37 %int_1
+         %40 = OpLoad %v2float %39
+         %41 = OpAccessChain %_ptr_Function_v2float %output %int_1
+               OpStore %41 %40
+         %42 = OpLoad %HSOut %output
+               OpReturnValue %42
+               OpFunctionEnd
+%PatchHS_struct_VertexOutput_vf4_vf21_3__ = OpFunction %HSConstantOut None %23
+      %patch = OpFunctionParameter %_ptr_Function__arr_VertexOutput_uint_3
+         %26 = OpLabel
+   %output_0 = OpVariable %_ptr_Function_HSConstantOut Function
+        %135 = OpAccessChain %_ptr_Function_v2float %patch %int_0 %int_1
+        %136 = OpLoad %v2float %135
+        %137 = OpCompositeConstruct %v2float %float_1 %float_1
+        %138 = OpFAdd %v2float %137 %136
+        %139 = OpCompositeExtract %float %138 0
+        %140 = OpAccessChain %_ptr_Function_float %output_0 %int_0 %int_0
+               OpStore %140 %139
+        %141 = OpAccessChain %_ptr_Function_v2float %patch %int_0 %int_1
+        %142 = OpLoad %v2float %141
+        %143 = OpCompositeConstruct %v2float %float_1 %float_1
+        %144 = OpFAdd %v2float %143 %142
+        %145 = OpCompositeExtract %float %144 0
+        %146 = OpAccessChain %_ptr_Function_float %output_0 %int_0 %int_1
+               OpStore %146 %145
+        %147 = OpAccessChain %_ptr_Function_v2float %patch %int_0 %int_1
+        %148 = OpLoad %v2float %147
+        %149 = OpCompositeConstruct %v2float %float_1 %float_1
+        %150 = OpFAdd %v2float %149 %148
+        %151 = OpCompositeExtract %float %150 0
+        %152 = OpAccessChain %_ptr_Function_float %output_0 %int_0 %int_2
+               OpStore %152 %151
+        %153 = OpAccessChain %_ptr_Function_v2float %patch %int_0 %int_1
+        %154 = OpLoad %v2float %153
+        %155 = OpCompositeConstruct %v2float %float_1 %float_1
+        %156 = OpFAdd %v2float %155 %154
+        %157 = OpCompositeExtract %float %156 0
+        %158 = OpAccessChain %_ptr_Function_float %output_0 %int_1
+               OpStore %158 %157
+        %159 = OpLoad %HSConstantOut %output_0
+               OpReturnValue %159
+               OpFunctionEnd

3rdparty/spirv-cross/shaders-msl/desktop-only/tesc/arrayed-output.desktop.sso.tesc

@@ -0,0 +1,27 @@
+#version 450
+
+layout(vertices = 4) out;
+layout(location = 0) patch out vec3 vPatch[2];
+layout(location = 2) out vec3 vVertex[];
+layout(location = 0) in vec3 vInput[];
+
+void main()
+{
+        vVertex[gl_InvocationID] =
+                vInput[gl_InvocationID] +
+                vInput[gl_InvocationID ^ 1];
+
+        barrier();
+
+        if (gl_InvocationID == 0)
+        {
+                vPatch[0] = vec3(10.0);
+                vPatch[1] = vec3(20.0);
+                gl_TessLevelOuter[0] = 1.0;
+                gl_TessLevelOuter[1] = 2.0;
+                gl_TessLevelOuter[2] = 3.0;
+                gl_TessLevelOuter[3] = 4.0;
+                gl_TessLevelInner[0] = 1.0;
+                gl_TessLevelInner[1] = 2.0;
+        }
+}

3rdparty/spirv-cross/shaders-msl/desktop-only/tesc/basic.desktop.sso.tesc

@@ -0,0 +1,32 @@
+#version 450
+layout(vertices = 1) out;
+
+in gl_PerVertex
+{
+   vec4 gl_Position;
+} gl_in[gl_MaxPatchVertices];
+
+out gl_PerVertex
+{
+   vec4 gl_Position;
+} gl_out[1];
+
+layout(location = 0) patch out vec3 vFoo;
+
+void set_position()
+{
+    gl_out[gl_InvocationID].gl_Position = gl_in[0].gl_Position + gl_in[1].gl_Position;
+}
+
+void main()
+{
+    gl_TessLevelInner[0] = 8.9;
+    gl_TessLevelInner[1] = 6.9;
+    gl_TessLevelOuter[0] = 8.9;
+    gl_TessLevelOuter[1] = 6.9;
+    gl_TessLevelOuter[2] = 3.9;
+    gl_TessLevelOuter[3] = 4.9;
+    vFoo = vec3(1.0);
+
+    set_position();
+}

3rdparty/spirv-cross/shaders-msl/desktop-only/tesc/struct-copy.desktop.sso.tesc

@@ -0,0 +1,22 @@
+#version 450
+
+struct Boo
+{
+        vec3 a;
+        vec3 b;
+};
+
+layout(vertices = 4) out;
+layout(location = 0) out Boo vVertex[];
+layout(location = 0) in Boo vInput[];
+
+void main()
+{
+        vVertex[gl_InvocationID] = vInput[gl_InvocationID];
+        gl_TessLevelOuter[0] = 1.0;
+        gl_TessLevelOuter[1] = 2.0;
+        gl_TessLevelOuter[2] = 3.0;
+        gl_TessLevelOuter[3] = 4.0;
+        gl_TessLevelInner[0] = 1.0;
+        gl_TessLevelInner[1] = 2.0;
+}

3rdparty/spirv-cross/shaders-msl/tesc/basic.tesc

@@ -0,0 +1,17 @@
+#version 310 es
+#extension GL_EXT_tessellation_shader : require
+
+layout(location = 0) patch out vec3 vFoo;
+
+layout(vertices = 1) out;
+
+void main()
+{
+    gl_TessLevelInner[0] = 8.9;
+    gl_TessLevelInner[1] = 6.9;
+    gl_TessLevelOuter[0] = 8.9;
+    gl_TessLevelOuter[1] = 6.9;
+    gl_TessLevelOuter[2] = 3.9;
+    gl_TessLevelOuter[3] = 4.9;
+    vFoo = vec3(1.0);
+}

3rdparty/spirv-cross/shaders-msl/tesc/water_tess.tesc

@@ -0,0 +1,115 @@
+#version 310 es
+#extension GL_EXT_tessellation_shader : require
+
+layout(vertices = 1) out;
+layout(location = 0) in vec2 vPatchPosBase[];
+
+layout(std140) uniform UBO
+{
+    vec4 uScale;
+    highp vec3 uCamPos;
+    vec2 uPatchSize;
+    vec2 uMaxTessLevel;
+    float uDistanceMod;
+    vec4 uFrustum[6];
+};
+
+layout(location = 1) patch out vec2 vOutPatchPosBase;
+layout(location = 2) patch out vec4 vPatchLods;
+
+float lod_factor(vec2 pos_)
+{
+    vec2 pos = pos_ * uScale.xy;
+    vec3 dist_to_cam = uCamPos - vec3(pos.x, 0.0, pos.y);
+    float level = log2((length(dist_to_cam) + 0.0001) * uDistanceMod);
+    return clamp(level, 0.0, uMaxTessLevel.x);
+}
+
+float tess_level(float lod)
+{
+    return uMaxTessLevel.y * exp2(-lod);
+}
+
+vec4 tess_level(vec4 lod)
+{
+    return uMaxTessLevel.y * exp2(-lod);
+}
+
+// Guard band for vertex displacement.
+#define GUARD_BAND 10.0
+bool frustum_cull(vec2 p0)
+{
+    vec2 min_xz = (p0 - GUARD_BAND) * uScale.xy;
+    vec2 max_xz = (p0 + uPatchSize + GUARD_BAND) * uScale.xy;
+
+    vec3 bb_min = vec3(min_xz.x, -GUARD_BAND, min_xz.y);
+    vec3 bb_max = vec3(max_xz.x, +GUARD_BAND, max_xz.y);
+    vec3 center = 0.5 * (bb_min + bb_max);
+    float radius = 0.5 * length(bb_max - bb_min);
+
+    vec3 f0 = vec3(
+        dot(uFrustum[0], vec4(center, 1.0)),
+        dot(uFrustum[1], vec4(center, 1.0)),
+        dot(uFrustum[2], vec4(center, 1.0)));
+
+    vec3 f1 = vec3(
+        dot(uFrustum[3], vec4(center, 1.0)),
+        dot(uFrustum[4], vec4(center, 1.0)),
+        dot(uFrustum[5], vec4(center, 1.0)));
+
+    return !(any(lessThanEqual(f0, vec3(-radius))) || any(lessThanEqual(f1, vec3(-radius))));
+}
+
+void compute_tess_levels(vec2 p0)
+{
+    vOutPatchPosBase = p0;
+
+    float l00 = lod_factor(p0 + vec2(-0.5, -0.5) * uPatchSize);
+    float l10 = lod_factor(p0 + vec2(+0.5, -0.5) * uPatchSize);
+    float l20 = lod_factor(p0 + vec2(+1.5, -0.5) * uPatchSize);
+    float l01 = lod_factor(p0 + vec2(-0.5, +0.5) * uPatchSize);
+    float l11 = lod_factor(p0 + vec2(+0.5, +0.5) * uPatchSize);
+    float l21 = lod_factor(p0 + vec2(+1.5, +0.5) * uPatchSize);
+    float l02 = lod_factor(p0 + vec2(-0.5, +1.5) * uPatchSize);
+    float l12 = lod_factor(p0 + vec2(+0.5, +1.5) * uPatchSize);
+    float l22 = lod_factor(p0 + vec2(+1.5, +1.5) * uPatchSize);
+
+    vec4 lods = vec4(
+        dot(vec4(l01, l11, l02, l12), vec4(0.25)),
+        dot(vec4(l00, l10, l01, l11), vec4(0.25)),
+        dot(vec4(l10, l20, l11, l21), vec4(0.25)),
+        dot(vec4(l11, l21, l12, l22), vec4(0.25)));
+
+    vPatchLods = lods;
+
+    vec4 outer_lods = min(lods.xyzw, lods.yzwx);
+    vec4 levels = tess_level(outer_lods);
+    gl_TessLevelOuter[0] = levels.x;
+    gl_TessLevelOuter[1] = levels.y;
+    gl_TessLevelOuter[2] = levels.z;
+    gl_TessLevelOuter[3] = levels.w;
+
+    float min_lod = min(min(lods.x, lods.y), min(lods.z, lods.w));
+    float inner = tess_level(min(min_lod, l11));
+    gl_TessLevelInner[0] = inner;
+    gl_TessLevelInner[1] = inner;
+}
+
+void main()
+{
+    vec2 p0 = vPatchPosBase[0];
+    if (!frustum_cull(p0))
+    {
+        gl_TessLevelOuter[0] = -1.0;
+        gl_TessLevelOuter[1] = -1.0;
+        gl_TessLevelOuter[2] = -1.0;
+        gl_TessLevelOuter[3] = -1.0;
+        gl_TessLevelInner[0] = -1.0;
+        gl_TessLevelInner[1] = -1.0;
+    }
+    else
+    {
+        compute_tess_levels(p0);
+    }
+}
+

3rdparty/spirv-cross/shaders/asm/comp/recompile-block-naming.asm.comp

@@ -0,0 +1,140 @@
+; SPIR-V
+; Version: 1.0
+; Generator: Khronos Glslang Reference Front End; 7
+; Bound: 97
+; Schema: 0
+               OpCapability Shader
+          %1 = OpExtInstImport "GLSL.std.450"
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpSource HLSL 500
+               OpName %main "main"
+               OpName %_main_ "@main("
+               OpName %a "a"
+               OpName %byteAddrTemp "byteAddrTemp"
+               OpName %MyFirstBuffer "MyFirstBuffer"
+               OpMemberName %MyFirstBuffer 0 "@data"
+               OpName %MyFirstBuffer_0 "MyFirstBuffer"
+               OpName %b "b"
+               OpName %byteAddrTemp_0 "byteAddrTemp"
+               OpName %MySecondBuffer "MySecondBuffer"
+               OpName %byteAddrTemp_1 "byteAddrTemp"
+               OpName %MyThirdBuffer "MyThirdBuffer"
+               OpDecorate %_runtimearr_uint ArrayStride 4
+               OpMemberDecorate %MyFirstBuffer 0 Offset 0
+               OpDecorate %MyFirstBuffer BufferBlock
+               OpDecorate %MyFirstBuffer_0 DescriptorSet 0
+               OpDecorate %MyFirstBuffer_0 Binding 0
+               OpDecorate %MySecondBuffer DescriptorSet 0
+               OpDecorate %MySecondBuffer Binding 0
+               OpDecorate %MyThirdBuffer DescriptorSet 0
+               OpDecorate %MyThirdBuffer Binding 0
+       %void = OpTypeVoid
+          %3 = OpTypeFunction %void
+       %uint = OpTypeInt 32 0
+     %v4uint = OpTypeVector %uint 4
+%_ptr_Function_v4uint = OpTypePointer Function %v4uint
+        %int = OpTypeInt 32 1
+%_ptr_Function_int = OpTypePointer Function %int
+      %int_0 = OpConstant %int 0
+      %int_2 = OpConstant %int 2
+%_runtimearr_uint = OpTypeRuntimeArray %uint
+%MyFirstBuffer = OpTypeStruct %_runtimearr_uint
+%_ptr_Uniform_MyFirstBuffer = OpTypePointer Uniform %MyFirstBuffer
+%MyFirstBuffer_0 = OpVariable %_ptr_Uniform_MyFirstBuffer Uniform
+%_ptr_Uniform_uint = OpTypePointer Uniform %uint
+      %int_1 = OpConstant %int 1
+      %int_3 = OpConstant %int 3
+      %int_4 = OpConstant %int 4
+%MySecondBuffer = OpVariable %_ptr_Uniform_MyFirstBuffer Uniform
+%MyThirdBuffer = OpVariable %_ptr_Uniform_MyFirstBuffer Uniform
+     %uint_0 = OpConstant %uint 0
+     %uint_1 = OpConstant %uint 1
+     %uint_2 = OpConstant %uint 2
+     %uint_3 = OpConstant %uint 3
+       %main = OpFunction %void None %3
+          %5 = OpLabel
+         %96 = OpFunctionCall %void %_main_
+               OpReturn
+               OpFunctionEnd
+     %_main_ = OpFunction %void None %3
+          %7 = OpLabel
+          %a = OpVariable %_ptr_Function_v4uint Function
+%byteAddrTemp = OpVariable %_ptr_Function_int Function
+          %b = OpVariable %_ptr_Function_v4uint Function
+%byteAddrTemp_0 = OpVariable %_ptr_Function_int Function
+%byteAddrTemp_1 = OpVariable %_ptr_Function_int Function
+         %17 = OpShiftRightArithmetic %int %int_0 %int_2
+               OpStore %byteAddrTemp %17
+         %22 = OpLoad %int %byteAddrTemp
+         %24 = OpAccessChain %_ptr_Uniform_uint %MyFirstBuffer_0 %int_0 %22
+         %25 = OpLoad %uint %24
+         %26 = OpLoad %int %byteAddrTemp
+         %28 = OpIAdd %int %26 %int_1
+         %29 = OpAccessChain %_ptr_Uniform_uint %MyFirstBuffer_0 %int_0 %28
+         %30 = OpLoad %uint %29
+         %31 = OpLoad %int %byteAddrTemp
+         %32 = OpIAdd %int %31 %int_2
+         %33 = OpAccessChain %_ptr_Uniform_uint %MyFirstBuffer_0 %int_0 %32
+         %34 = OpLoad %uint %33
+         %35 = OpLoad %int %byteAddrTemp
+         %37 = OpIAdd %int %35 %int_3
+         %38 = OpAccessChain %_ptr_Uniform_uint %MyFirstBuffer_0 %int_0 %37
+         %39 = OpLoad %uint %38
+         %40 = OpCompositeConstruct %v4uint %25 %30 %34 %39
+               OpStore %a %40
+         %44 = OpShiftRightArithmetic %int %int_4 %int_2
+               OpStore %byteAddrTemp_0 %44
+         %46 = OpLoad %int %byteAddrTemp_0
+         %47 = OpAccessChain %_ptr_Uniform_uint %MySecondBuffer %int_0 %46
+         %48 = OpLoad %uint %47
+         %49 = OpLoad %int %byteAddrTemp_0
+         %50 = OpIAdd %int %49 %int_1
+         %51 = OpAccessChain %_ptr_Uniform_uint %MySecondBuffer %int_0 %50
+         %52 = OpLoad %uint %51
+         %53 = OpLoad %int %byteAddrTemp_0
+         %54 = OpIAdd %int %53 %int_2
+         %55 = OpAccessChain %_ptr_Uniform_uint %MySecondBuffer %int_0 %54
+         %56 = OpLoad %uint %55
+         %57 = OpLoad %int %byteAddrTemp_0
+         %58 = OpIAdd %int %57 %int_3
+         %59 = OpAccessChain %_ptr_Uniform_uint %MySecondBuffer %int_0 %58
+         %60 = OpLoad %uint %59
+         %61 = OpCompositeConstruct %v4uint %48 %52 %56 %60
+               OpStore %b %61
+         %63 = OpShiftRightArithmetic %int %int_0 %int_2
+               OpStore %byteAddrTemp_1 %63
+         %65 = OpLoad %int %byteAddrTemp_1
+         %66 = OpLoad %v4uint %a
+         %67 = OpLoad %v4uint %b
+         %68 = OpIAdd %v4uint %66 %67
+         %70 = OpCompositeExtract %uint %68 0
+         %71 = OpAccessChain %_ptr_Uniform_uint %MyThirdBuffer %int_0 %65
+               OpStore %71 %70
+         %72 = OpLoad %int %byteAddrTemp_1
+         %73 = OpIAdd %int %72 %int_1
+         %74 = OpLoad %v4uint %a
+         %75 = OpLoad %v4uint %b
+         %76 = OpIAdd %v4uint %74 %75
+         %78 = OpCompositeExtract %uint %76 1
+         %79 = OpAccessChain %_ptr_Uniform_uint %MyThirdBuffer %int_0 %73
+               OpStore %79 %78
+         %80 = OpLoad %int %byteAddrTemp_1
+         %81 = OpIAdd %int %80 %int_2
+         %82 = OpLoad %v4uint %a
+         %83 = OpLoad %v4uint %b
+         %84 = OpIAdd %v4uint %82 %83
+         %86 = OpCompositeExtract %uint %84 2
+         %87 = OpAccessChain %_ptr_Uniform_uint %MyThirdBuffer %int_0 %81
+               OpStore %87 %86
+         %88 = OpLoad %int %byteAddrTemp_1
+         %89 = OpIAdd %int %88 %int_3
+         %90 = OpLoad %v4uint %a
+         %91 = OpLoad %v4uint %b
+         %92 = OpIAdd %v4uint %90 %91
+         %94 = OpCompositeExtract %uint %92 3
+         %95 = OpAccessChain %_ptr_Uniform_uint %MyThirdBuffer %int_0 %89
+               OpStore %95 %94
+               OpReturn
+               OpFunctionEnd

3rdparty/spirv-cross/spirv_common.hpp

@@ -1381,6 +1381,8 @@ struct Meta
 		{
 			uint32_t packed_type = 0;
 			bool packed = false;
+			uint32_t ib_member_index = -1;
+			uint32_t ib_orig_id = 0;
 		} extended;
 	};
 

3rdparty/spirv-cross/spirv_cross.cpp

@@ -1094,6 +1094,22 @@ const SPIRType &Compiler::get_variable_data_type(const SPIRVariable &var) const
 	return get<SPIRType>(get_variable_data_type_id(var));
 }
 
+SPIRType &Compiler::get_variable_element_type(const SPIRVariable &var)
+{
+	SPIRType *type = &get_variable_data_type(var);
+	if (is_array(*type))
+		type = &get<SPIRType>(type->parent_type);
+	return *type;
+}
+
+const SPIRType &Compiler::get_variable_element_type(const SPIRVariable &var) const
+{
+	const SPIRType *type = &get_variable_data_type(var);
+	if (is_array(*type))
+		type = &get<SPIRType>(type->parent_type);
+	return *type;
+}
+
 bool Compiler::is_sampled_image_type(const SPIRType &type)
 {
 	return (type.basetype == SPIRType::Image || type.basetype == SPIRType::SampledImage) && type.image.sampled == 1 &&
@@ -1183,6 +1199,14 @@ void Compiler::set_extended_decoration(uint32_t id, ExtendedDecorations decorati
 	case SPIRVCrossDecorationPackedType:
 		dec.extended.packed_type = value;
 		break;
+
+	case SPIRVCrossDecorationInterfaceMemberIndex:
+		dec.extended.ib_member_index = value;
+		break;
+
+	case SPIRVCrossDecorationInterfaceOrigID:
+		dec.extended.ib_orig_id = value;
+		break;
 	}
 }
 
@@ -1201,6 +1225,14 @@ void Compiler::set_extended_member_decoration(uint32_t type, uint32_t index, Ext
 	case SPIRVCrossDecorationPackedType:
 		dec.extended.packed_type = value;
 		break;
+
+	case SPIRVCrossDecorationInterfaceMemberIndex:
+		dec.extended.ib_member_index = value;
+		break;
+
+	case SPIRVCrossDecorationInterfaceOrigID:
+		dec.extended.ib_orig_id = value;
+		break;
 	}
 }
 
@@ -1218,6 +1250,12 @@ uint32_t Compiler::get_extended_decoration(uint32_t id, ExtendedDecorations deco
 
 	case SPIRVCrossDecorationPackedType:
 		return dec.extended.packed_type;
+
+	case SPIRVCrossDecorationInterfaceMemberIndex:
+		return dec.extended.ib_member_index;
+
+	case SPIRVCrossDecorationInterfaceOrigID:
+		return dec.extended.ib_orig_id;
 	}
 
 	return 0;
@@ -1240,6 +1278,12 @@ uint32_t Compiler::get_extended_member_decoration(uint32_t type, uint32_t index,
 
 	case SPIRVCrossDecorationPackedType:
 		return dec.extended.packed_type;
+
+	case SPIRVCrossDecorationInterfaceMemberIndex:
+		return dec.extended.ib_member_index;
+
+	case SPIRVCrossDecorationInterfaceOrigID:
+		return dec.extended.ib_orig_id;
 	}
 
 	return 0;
@@ -1259,6 +1303,12 @@ bool Compiler::has_extended_decoration(uint32_t id, ExtendedDecorations decorati
 
 	case SPIRVCrossDecorationPackedType:
 		return dec.extended.packed_type != 0;
+
+	case SPIRVCrossDecorationInterfaceMemberIndex:
+		return dec.extended.ib_member_index != uint32_t(-1);
+
+	case SPIRVCrossDecorationInterfaceOrigID:
+		return dec.extended.ib_orig_id != 0;
 	}
 
 	return false;
@@ -1281,6 +1331,12 @@ bool Compiler::has_extended_member_decoration(uint32_t type, uint32_t index, Ext
 
 	case SPIRVCrossDecorationPackedType:
 		return dec.extended.packed_type != 0;
+
+	case SPIRVCrossDecorationInterfaceMemberIndex:
+		return dec.extended.ib_member_index != uint32_t(-1);
+
+	case SPIRVCrossDecorationInterfaceOrigID:
+		return dec.extended.ib_orig_id != 0;
 	}
 
 	return false;
@@ -1298,6 +1354,14 @@ void Compiler::unset_extended_decoration(uint32_t id, ExtendedDecorations decora
 	case SPIRVCrossDecorationPackedType:
 		dec.extended.packed_type = 0;
 		break;
+
+	case SPIRVCrossDecorationInterfaceMemberIndex:
+		dec.extended.ib_member_index = -1;
+		break;
+
+	case SPIRVCrossDecorationInterfaceOrigID:
+		dec.extended.ib_orig_id = 0;
+		break;
 	}
 }
 
@@ -1315,6 +1379,14 @@ void Compiler::unset_extended_member_decoration(uint32_t type, uint32_t index, E
 	case SPIRVCrossDecorationPackedType:
 		dec.extended.packed_type = 0;
 		break;
+
+	case SPIRVCrossDecorationInterfaceMemberIndex:
+		dec.extended.ib_member_index = -1;
+		break;
+
+	case SPIRVCrossDecorationInterfaceOrigID:
+		dec.extended.ib_orig_id = 0;
+		break;
 	}
 }
 
@@ -3527,7 +3599,7 @@ bool Compiler::ActiveBuiltinHandler::handle(spv::Op opcode, const uint32_t *args
 		auto *type = &compiler.get_variable_data_type(*var);
 
 		auto &flags =
-		    type->storage == StorageClassInput ? compiler.active_input_builtins : compiler.active_output_builtins;
+		    var->storage == StorageClassInput ? compiler.active_input_builtins : compiler.active_output_builtins;
 
 		uint32_t count = length - 3;
 		args += 3;

3rdparty/spirv-cross/spirv_cross.hpp

@@ -117,7 +117,9 @@ struct EntryPoint
 enum ExtendedDecorations
 {
 	SPIRVCrossDecorationPacked,
-	SPIRVCrossDecorationPackedType
+	SPIRVCrossDecorationPackedType,
+	SPIRVCrossDecorationInterfaceMemberIndex,
+	SPIRVCrossDecorationInterfaceOrigID,
 };
 
 class Compiler
@@ -183,27 +185,6 @@ public:
 	// Gets the SPIR-V type of a variable.
 	const SPIRType &get_type_from_variable(uint32_t id) const;
 
-	// Gets the id of SPIR-V type underlying the given type_id, which might be a pointer.
-	uint32_t get_pointee_type_id(uint32_t type_id) const;
-
-	// Gets the SPIR-V type underlying the given type, which might be a pointer.
-	const SPIRType &get_pointee_type(const SPIRType &type) const;
-
-	// Gets the SPIR-V type underlying the given type_id, which might be a pointer.
-	const SPIRType &get_pointee_type(uint32_t type_id) const;
-
-	// Gets the ID of the SPIR-V type underlying a variable.
-	uint32_t get_variable_data_type_id(const SPIRVariable &var) const;
-
-	// Gets the SPIR-V type underlying a variable.
-	SPIRType &get_variable_data_type(const SPIRVariable &var);
-
-	// Gets the SPIR-V type underlying a variable.
-	const SPIRType &get_variable_data_type(const SPIRVariable &var) const;
-
-	// Returns if the given type refers to a sampled image.
-	bool is_sampled_image_type(const SPIRType &type);
-
 	// Gets the underlying storage class for an OpVariable.
 	spv::StorageClass get_storage_class(uint32_t id) const;
 
@@ -603,6 +584,33 @@ protected:
 			return nullptr;
 	}
 
+	// Gets the id of SPIR-V type underlying the given type_id, which might be a pointer.
+	uint32_t get_pointee_type_id(uint32_t type_id) const;
+
+	// Gets the SPIR-V type underlying the given type, which might be a pointer.
+	const SPIRType &get_pointee_type(const SPIRType &type) const;
+
+	// Gets the SPIR-V type underlying the given type_id, which might be a pointer.
+	const SPIRType &get_pointee_type(uint32_t type_id) const;
+
+	// Gets the ID of the SPIR-V type underlying a variable.
+	uint32_t get_variable_data_type_id(const SPIRVariable &var) const;
+
+	// Gets the SPIR-V type underlying a variable.
+	SPIRType &get_variable_data_type(const SPIRVariable &var);
+
+	// Gets the SPIR-V type underlying a variable.
+	const SPIRType &get_variable_data_type(const SPIRVariable &var) const;
+
+	// Gets the SPIR-V element type underlying an array variable.
+	SPIRType &get_variable_element_type(const SPIRVariable &var);
+
+	// Gets the SPIR-V element type underlying an array variable.
+	const SPIRType &get_variable_element_type(const SPIRVariable &var) const;
+
+	// Returns if the given type refers to a sampled image.
+	bool is_sampled_image_type(const SPIRType &type);
+
 	const SPIREntryPoint &get_entry_point() const;
 	SPIREntryPoint &get_entry_point();
 

3rdparty/spirv-cross/spirv_glsl.cpp

@@ -262,13 +262,7 @@ void CompilerGLSL::reset()
 	expression_usage_counts.clear();
 	forwarded_temporaries.clear();
 
-	resource_names.clear();
-	block_input_names.clear();
-	block_output_names.clear();
-	block_ubo_names.clear();
-	block_ssbo_names.clear();
-	block_names.clear();
-	function_overloads.clear();
+	reset_name_caches();
 
 	ir.for_each_typed_id<SPIRFunction>([&](uint32_t, SPIRFunction &func) {
 		func.active = false;
@@ -1561,6 +1555,10 @@ void CompilerGLSL::emit_buffer_block_native(const SPIRVariable &var)
 		i++;
 	}
 
+	// var.self can be used as a backup name for the block name,
+	// so we need to make sure we don't disturb the name here on a recompile.
+	// It will need to be reset if we have to recompile.
+	preserve_alias_on_reset(var.self);
 	add_resource_name(var.self);
 	end_scope_decl(to_name(var.self) + type_to_array_glsl(type));
 	statement("");
@@ -10984,3 +10982,23 @@ void CompilerGLSL::bitcast_to_builtin_store(uint32_t target_id, std::string &exp
 void CompilerGLSL::emit_block_hints(const SPIRBlock &)
 {
 }
+
+void CompilerGLSL::preserve_alias_on_reset(uint32_t id)
+{
+	preserved_aliases[id] = get_name(id);
+}
+
+void CompilerGLSL::reset_name_caches()
+{
+	for (auto &preserved : preserved_aliases)
+		set_name(preserved.first, preserved.second);
+
+	preserved_aliases.clear();
+	resource_names.clear();
+	block_input_names.clear();
+	block_output_names.clear();
+	block_ubo_names.clear();
+	block_ssbo_names.clear();
+	block_names.clear();
+	function_overloads.clear();
+}

3rdparty/spirv-cross/spirv_glsl.hpp

@@ -358,6 +358,9 @@ protected:
 	std::unordered_set<std::string> block_ssbo_names;
 	std::unordered_set<std::string> block_names; // A union of all block_*_names.
 	std::unordered_map<std::string, std::unordered_set<uint64_t>> function_overloads;
+	std::unordered_map<uint32_t, std::string> preserved_aliases;
+	void preserve_alias_on_reset(uint32_t id);
+	void reset_name_caches();
 
 	bool processing_entry_point = false;
 

3rdparty/spirv-cross/spirv_hlsl.cpp

@@ -1838,6 +1838,10 @@ void CompilerHLSL::emit_buffer_block(const SPIRVariable &var)
 			declared_block_names[var.self] = buffer_name;
 
 			type.member_name_cache.clear();
+			// var.self can be used as a backup name for the block name,
+			// so we need to make sure we don't disturb the name here on a recompile.
+			// It will need to be reset if we have to recompile.
+			preserve_alias_on_reset(var.self);
 			add_resource_name(var.self);
 			statement("cbuffer ", buffer_name, to_resource_binding(var));
 			begin_scope();

3rdparty/spirv-cross/spirv_msl.hpp

@@ -170,6 +170,9 @@ public:
 		uint32_t aux_buffer_index = 30;
 		uint32_t indirect_params_buffer_index = 29;
 		uint32_t shader_output_buffer_index = 28;
+		uint32_t shader_patch_output_buffer_index = 27;
+		uint32_t shader_tess_factor_buffer_index = 26;
+		uint32_t shader_input_wg_index = 0;
 		bool enable_point_size_builtin = true;
 		bool disable_rasterization = false;
 		bool capture_output_to_buffer = false;
@@ -231,7 +234,8 @@ public:
 	// rasterization if vertex shader requires rasterization to be disabled.
 	bool get_is_rasterization_disabled() const
 	{
-		return is_rasterization_disabled && (get_entry_point().model == spv::ExecutionModelVertex);
+		return is_rasterization_disabled && (get_entry_point().model == spv::ExecutionModelVertex ||
+		                                     get_entry_point().model == spv::ExecutionModelTessellationControl);
 	}
 
 	// Provide feedback to calling API to allow it to pass an auxiliary
@@ -248,6 +252,20 @@ public:
 		return capture_output_to_buffer && stage_out_var_id != 0;
 	}
 
+	// Provide feedback to calling API to allow it to pass a patch output
+	// buffer if the shader needs it.
+	bool needs_patch_output_buffer() const
+	{
+		return capture_output_to_buffer && patch_stage_out_var_id != 0;
+	}
+
+	// Provide feedback to calling API to allow it to pass an input threadgroup
+	// buffer if the shader needs it.
+	bool needs_input_threadgroup_mem() const
+	{
+		return capture_output_to_buffer && stage_in_var_id != 0;
+	}
+
 	// An enum of SPIR-V functions that are implemented in additional
 	// source code that is added to the shader if necessary.
 	enum SPVFuncImpl
@@ -384,19 +402,24 @@ protected:
 	void extract_global_variables_from_function(uint32_t func_id, std::set<uint32_t> &added_arg_ids,
 	                                            std::unordered_set<uint32_t> &global_var_ids,
 	                                            std::unordered_set<uint32_t> &processed_func_ids);
-	uint32_t add_interface_block(spv::StorageClass storage);
+	uint32_t add_interface_block(spv::StorageClass storage, bool patch = false);
+	uint32_t add_interface_block_pointer(uint32_t ib_var_id, spv::StorageClass storage);
 
 	void add_variable_to_interface_block(spv::StorageClass storage, const std::string &ib_var_ref, SPIRType &ib_type,
-	                                     SPIRVariable &var);
+	                                     SPIRVariable &var, bool strip_array);
 	void add_composite_variable_to_interface_block(spv::StorageClass storage, const std::string &ib_var_ref,
-	                                               SPIRType &ib_type, SPIRVariable &var);
+	                                               SPIRType &ib_type, SPIRVariable &var, bool strip_array);
 	void add_plain_variable_to_interface_block(spv::StorageClass storage, const std::string &ib_var_ref,
-	                                           SPIRType &ib_type, SPIRVariable &var);
+	                                           SPIRType &ib_type, SPIRVariable &var, bool strip_array);
 	void add_plain_member_variable_to_interface_block(spv::StorageClass storage, const std::string &ib_var_ref,
-	                                                  SPIRType &ib_type, SPIRVariable &var, uint32_t index);
+	                                                  SPIRType &ib_type, SPIRVariable &var, uint32_t index,
+	                                                  bool strip_array);
 	void add_composite_member_variable_to_interface_block(spv::StorageClass storage, const std::string &ib_var_ref,
-	                                                      SPIRType &ib_type, SPIRVariable &var, uint32_t index);
-	uint32_t get_accumulated_member_location(const SPIRVariable &var, uint32_t mbr_idx);
+	                                                      SPIRType &ib_type, SPIRVariable &var, uint32_t index,
+	                                                      bool strip_array);
+	uint32_t get_accumulated_member_location(const SPIRVariable &var, uint32_t mbr_idx, bool strip_array);
+
+	void fix_up_interface_member_indices(spv::StorageClass storage, uint32_t ib_type_id);
 
 	void mark_location_as_used_by_shader(uint32_t location, spv::StorageClass storage);
 	uint32_t ensure_correct_builtin_type(uint32_t type_id, spv::BuiltIn builtin);
@@ -431,7 +454,10 @@ protected:
 	MSLStructMemberKey get_struct_member_key(uint32_t type_id, uint32_t index);
 	std::string get_argument_address_space(const SPIRVariable &argument);
 	std::string get_type_address_space(const SPIRType &type);
+	SPIRType &get_stage_in_struct_type();
 	SPIRType &get_stage_out_struct_type();
+	SPIRType &get_patch_stage_out_struct_type();
+	std::string get_tess_factor_struct_name();
 	void emit_atomic_func_op(uint32_t result_type, uint32_t result_id, const char *op, uint32_t mem_order_1,
 	                         uint32_t mem_order_2, bool has_mem_order_2, uint32_t op0, uint32_t op1 = 0,
 	                         bool op1_is_pointer = false, bool op1_is_literal = false, uint32_t op2 = 0);
@@ -448,6 +474,8 @@ protected:
 	uint32_t builtin_base_vertex_id = 0;
 	uint32_t builtin_instance_idx_id = 0;
 	uint32_t builtin_base_instance_id = 0;
+	uint32_t builtin_invocation_id_id = 0;
+	uint32_t builtin_primitive_id_id = 0;
 	uint32_t aux_buffer_id = 0;
 
 	void bitcast_to_builtin_store(uint32_t target_id, std::string &expr, const SPIRType &expr_type) override;
@@ -456,6 +484,8 @@ protected:
 
 	void analyze_sampled_image_usage();
 
+	bool emit_tessellation_control_access_chain(const uint32_t *ops, uint32_t length);
+
 	Options msl_options;
 	std::set<SPVFuncImpl> spv_function_implementations;
 	std::unordered_map<uint32_t, MSLVertexAttr *> vtx_attrs_by_location;
@@ -468,6 +498,9 @@ protected:
 	MSLResourceBinding next_metal_resource_index;
 	uint32_t stage_in_var_id = 0;
 	uint32_t stage_out_var_id = 0;
+	uint32_t patch_stage_out_var_id = 0;
+	uint32_t stage_in_ptr_var_id = 0;
+	uint32_t stage_out_ptr_var_id = 0;
 	bool has_sampled_images = false;
 	bool needs_vertex_idx_arg = false;
 	bool needs_instance_idx_arg = false;
@@ -478,9 +511,13 @@ protected:
 	std::string qual_pos_var_name;
 	std::string stage_in_var_name = "in";
 	std::string stage_out_var_name = "out";
+	std::string patch_stage_out_var_name = "patchOut";
 	std::string sampler_name_suffix = "Smplr";
 	std::string swizzle_name_suffix = "Swzl";
+	std::string input_wg_var_name = "gl_in";
 	std::string output_buffer_var_name = "spvOut";
+	std::string patch_output_buffer_var_name = "spvPatchOut";
+	std::string tess_factor_buffer_var_name = "spvTessLevel";
 	spv::Op previous_instruction_opcode = spv::OpNop;
 
 	std::unordered_map<uint32_t, MSLConstexprSampler> constexpr_samplers;

3rdparty/spirv-tools/Android.mk

@@ -99,6 +99,7 @@ SPVTOOLS_OPT_SRC_FILES := \
 		source/opt/eliminate_dead_constant_pass.cpp \
 		source/opt/eliminate_dead_functions_pass.cpp \
 		source/opt/eliminate_dead_functions_util.cpp \
+		source/opt/eliminate_dead_members_pass.cpp \
 		source/opt/feature_manager.cpp \
 		source/opt/flatten_decoration_pass.cpp \
 		source/opt/fold.cpp \

3rdparty/spirv-tools/BUILD.gn

@@ -305,7 +305,6 @@ source_set("spvtools_headers") {
 
 static_library("spvtools") {
   deps = [
-    ":spvtools_core_enums_unified1",
     ":spvtools_core_tables_unified1",
     ":spvtools_generators_inc",
     ":spvtools_glsl_tables_glsl1-0",
@@ -376,6 +375,7 @@ static_library("spvtools") {
   ]
 
   public_deps = [
+    ":spvtools_core_enums_unified1",
     ":spvtools_headers",
   ]
 
@@ -498,6 +498,8 @@ static_library("spvtools_opt") {
     "source/opt/eliminate_dead_functions_pass.h",
     "source/opt/eliminate_dead_functions_util.cpp",
     "source/opt/eliminate_dead_functions_util.h",
+    "source/opt/eliminate_dead_members_pass.cpp",
+    "source/opt/eliminate_dead_members_pass.h",
     "source/opt/feature_manager.cpp",
     "source/opt/feature_manager.h",
     "source/opt/flatten_decoration_pass.cpp",

3rdparty/spirv-tools/include/generated/build-version.inc

@@ -1 +1 @@
-"v2019.2-dev", "SPIRV-Tools v2019.2-dev c3405c13ff287fcc881a3244dafd761d40edaa50"
+"v2019.2-dev", "SPIRV-Tools v2019.2-dev f237f9cff729a0b78c8e979a4da7f4428bfeeaf2"

3rdparty/spirv-tools/include/spirv-tools/optimizer.hpp

@@ -226,6 +226,11 @@ Optimizer::PassToken CreateStripReflectInfoPass();
 // functions are not needed because they will never be called.
 Optimizer::PassToken CreateEliminateDeadFunctionsPass();
 
+// Creates an eliminate-dead-members pass.
+// An eliminate-dead-members pass will remove all unused members of structures.
+// This will not affect the data layout of the remaining members.
+Optimizer::PassToken CreateEliminateDeadMembersPass();
+
 // Creates a set-spec-constant-default-value pass from a mapping from spec-ids
 // to the default values in the form of string.
 // A set-spec-constant-default-value pass sets the default values for the

3rdparty/spirv-tools/source/opt/CMakeLists.txt

@@ -38,6 +38,7 @@ set(SPIRV_TOOLS_OPT_SOURCES
   eliminate_dead_constant_pass.h
   eliminate_dead_functions_pass.h
   eliminate_dead_functions_util.h
+  eliminate_dead_members_pass.h
   feature_manager.h
   flatten_decoration_pass.h
   fold.h
@@ -133,6 +134,7 @@ set(SPIRV_TOOLS_OPT_SOURCES
   eliminate_dead_constant_pass.cpp
   eliminate_dead_functions_pass.cpp
   eliminate_dead_functions_util.cpp
+  eliminate_dead_members_pass.cpp
   feature_manager.cpp
   flatten_decoration_pass.cpp
   fold.cpp

3rdparty/spirv-tools/source/opt/block_merge_util.cpp

@@ -86,6 +86,11 @@ bool CanMergeWithSuccessor(IRContext* context, BasicBlock* block) {
     return false;
   }
 
+  // Don't bother trying to merge unreachable blocks.
+  if (auto dominators = context->GetDominatorAnalysis(block->GetParent())) {
+    if (!dominators->IsReachable(block)) return false;
+  }
+
   Instruction* merge_inst = block->GetMergeInst();
   const bool pred_is_header = IsHeader(block);
   if (pred_is_header && lab_id != merge_inst->GetSingleWordInOperand(0u)) {

3rdparty/spirv-tools/source/opt/eliminate_dead_members_pass.cpp

@@ -0,0 +1,624 @@
+// Copyright (c) 2019 Google LLC
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+#include "source/opt/eliminate_dead_members_pass.h"
+
+#include "ir_builder.h"
+#include "source/opt/ir_context.h"
+
+namespace {
+const uint32_t kRemovedMember = 0xFFFFFFFF;
+}
+
+namespace spvtools {
+namespace opt {
+
+Pass::Status EliminateDeadMembersPass::Process() {
+  if (!context()->get_feature_mgr()->HasCapability(SpvCapabilityShader))
+    return Status::SuccessWithoutChange;
+
+  FindLiveMembers();
+  if (RemoveDeadMembers()) {
+    return Status::SuccessWithChange;
+  }
+  return Status::SuccessWithoutChange;
+}
+
+void EliminateDeadMembersPass::FindLiveMembers() {
+  // Until we have implemented the rewritting of OpSpecConsantOp instructions,
+  // we have to mark them as fully used just to be safe.
+  for (auto& inst : get_module()->types_values()) {
+    if (inst.opcode() != SpvOpSpecConstantOp) {
+      continue;
+    }
+    MarkTypeAsFullyUsed(inst.type_id());
+  }
+
+  for (const Function& func : *get_module()) {
+    FindLiveMembers(func);
+  }
+}
+
+void EliminateDeadMembersPass::FindLiveMembers(const Function& function) {
+  function.ForEachInst(
+      [this](const Instruction* inst) { FindLiveMembers(inst); });
+}
+
+void EliminateDeadMembersPass::FindLiveMembers(const Instruction* inst) {
+  switch (inst->opcode()) {
+    case SpvOpStore:
+      MarkMembersAsLiveForStore(inst);
+      break;
+    case SpvOpCopyMemory:
+    case SpvOpCopyMemorySized:
+      MarkMembersAsLiveForCopyMemory(inst);
+      break;
+    case SpvOpCompositeExtract:
+      MarkMembersAsLiveForExtract(inst);
+      break;
+    case SpvOpAccessChain:
+    case SpvOpInBoundsAccessChain:
+    case SpvOpPtrAccessChain:
+    case SpvOpInBoundsPtrAccessChain:
+      MarkMembersAsLiveForAccessChain(inst);
+      break;
+    case SpvOpReturnValue:
+      // This should be an issue only if we are returning from the entry point.
+      // However, for now I will keep it more conservative because functions are
+      // often inlined leaving only the entry points.
+      MarkOperandTypeAsFullyUsed(inst, 0);
+      break;
+    case SpvOpArrayLength:
+      MarkMembersAsLiveForArrayLength(inst);
+      break;
+    case SpvOpLoad:
+    case SpvOpCompositeInsert:
+    case SpvOpCompositeConstruct:
+      break;
+    default:
+      // This path is here for safety.  All instructions that can reference
+      // structs in a function body should be handled above.  However, this will
+      // keep the pass valid, but not optimal, as new instructions get added
+      // or if something was missed.
+      MarkStructOperandsAsFullyUsed(inst);
+      break;
+  }
+}
+
+void EliminateDeadMembersPass::MarkMembersAsLiveForStore(
+    const Instruction* inst) {
+  // We should only have to mark the members as live if the store is to
+  // memory that is read outside of the shader.  Other passes can remove all
+  // store to memory that is not visible outside of the shader, so we do not
+  // complicate the code for now.
+  assert(inst->opcode() == SpvOpStore);
+  uint32_t object_id = inst->GetSingleWordInOperand(1);
+  Instruction* object_inst = context()->get_def_use_mgr()->GetDef(object_id);
+  uint32_t object_type_id = object_inst->type_id();
+  MarkTypeAsFullyUsed(object_type_id);
+}
+
+void EliminateDeadMembersPass::MarkTypeAsFullyUsed(uint32_t type_id) {
+  Instruction* type_inst = get_def_use_mgr()->GetDef(type_id);
+  assert(type_inst != nullptr);
+  if (type_inst->opcode() != SpvOpTypeStruct) {
+    return;
+  }
+
+  // Mark every member of the current struct as used.
+  for (uint32_t i = 0; i < type_inst->NumInOperands(); ++i) {
+    used_members_[type_id].insert(i);
+  }
+
+  // Mark any sub struct as fully used.
+  for (uint32_t i = 0; i < type_inst->NumInOperands(); ++i) {
+    MarkTypeAsFullyUsed(type_inst->GetSingleWordInOperand(i));
+  }
+}
+
+void EliminateDeadMembersPass::MarkMembersAsLiveForCopyMemory(
+    const Instruction* inst) {
+  uint32_t target_id = inst->GetSingleWordInOperand(0);
+  Instruction* target_inst = get_def_use_mgr()->GetDef(target_id);
+  uint32_t pointer_type_id = target_inst->type_id();
+  Instruction* pointer_type_inst = get_def_use_mgr()->GetDef(pointer_type_id);
+  uint32_t type_id = pointer_type_inst->GetSingleWordInOperand(1);
+  MarkTypeAsFullyUsed(type_id);
+}
+
+void EliminateDeadMembersPass::MarkMembersAsLiveForExtract(
+    const Instruction* inst) {
+  assert(inst->opcode() == SpvOpCompositeExtract);
+
+  uint32_t composite_id = inst->GetSingleWordInOperand(0);
+  Instruction* composite_inst = get_def_use_mgr()->GetDef(composite_id);
+  uint32_t type_id = composite_inst->type_id();
+
+  for (uint32_t i = 1; i < inst->NumInOperands(); ++i) {
+    Instruction* type_inst = get_def_use_mgr()->GetDef(type_id);
+    uint32_t member_idx = inst->GetSingleWordInOperand(i);
+    switch (type_inst->opcode()) {
+      case SpvOpTypeStruct:
+        used_members_[type_id].insert(member_idx);
+        type_id = type_inst->GetSingleWordInOperand(member_idx);
+        break;
+      case SpvOpTypeArray:
+      case SpvOpTypeRuntimeArray:
+      case SpvOpTypeVector:
+      case SpvOpTypeMatrix:
+        type_id = type_inst->GetSingleWordInOperand(0);
+        break;
+      default:
+        assert(false);
+    }
+  }
+}
+
+void EliminateDeadMembersPass::MarkMembersAsLiveForAccessChain(
+    const Instruction* inst) {
+  assert(inst->opcode() == SpvOpAccessChain ||
+         inst->opcode() == SpvOpInBoundsAccessChain ||
+         inst->opcode() == SpvOpPtrAccessChain ||
+         inst->opcode() == SpvOpInBoundsPtrAccessChain);
+
+  uint32_t pointer_id = inst->GetSingleWordInOperand(0);
+  Instruction* pointer_inst = get_def_use_mgr()->GetDef(pointer_id);
+  uint32_t pointer_type_id = pointer_inst->type_id();
+  Instruction* pointer_type_inst = get_def_use_mgr()->GetDef(pointer_type_id);
+  uint32_t type_id = pointer_type_inst->GetSingleWordInOperand(1);
+
+  analysis::ConstantManager* const_mgr = context()->get_constant_mgr();
+
+  // For a pointer access chain, we need to skip the |element| index.  It is not
+  // a reference to the member of a struct, and it does not change the type.
+  uint32_t i = (inst->opcode() == SpvOpAccessChain ||
+                        inst->opcode() == SpvOpInBoundsAccessChain
+                    ? 1
+                    : 2);
+  for (; i < inst->NumInOperands(); ++i) {
+    Instruction* type_inst = get_def_use_mgr()->GetDef(type_id);
+    switch (type_inst->opcode()) {
+      case SpvOpTypeStruct: {
+        const analysis::IntConstant* member_idx =
+            const_mgr->FindDeclaredConstant(inst->GetSingleWordInOperand(i))
+                ->AsIntConstant();
+        assert(member_idx);
+        if (member_idx->type()->AsInteger()->width() == 32) {
+          used_members_[type_id].insert(member_idx->GetU32());
+          type_id = type_inst->GetSingleWordInOperand(member_idx->GetU32());
+        } else {
+          used_members_[type_id].insert(
+              static_cast<uint32_t>(member_idx->GetU64()));
+          type_id = type_inst->GetSingleWordInOperand(
+              static_cast<uint32_t>(member_idx->GetU64()));
+        }
+      } break;
+      case SpvOpTypeArray:
+      case SpvOpTypeRuntimeArray:
+      case SpvOpTypeVector:
+      case SpvOpTypeMatrix:
+        type_id = type_inst->GetSingleWordInOperand(0);
+        break;
+      default:
+        assert(false);
+    }
+  }
+}
+
+void EliminateDeadMembersPass::MarkOperandTypeAsFullyUsed(
+    const Instruction* inst, uint32_t in_idx) {
+  uint32_t op_id = inst->GetSingleWordInOperand(in_idx);
+  Instruction* op_inst = get_def_use_mgr()->GetDef(op_id);
+  MarkTypeAsFullyUsed(op_inst->type_id());
+}
+
+void EliminateDeadMembersPass::MarkMembersAsLiveForArrayLength(
+    const Instruction* inst) {
+  assert(inst->opcode() == SpvOpArrayLength);
+  uint32_t object_id = inst->GetSingleWordInOperand(0);
+  Instruction* object_inst = get_def_use_mgr()->GetDef(object_id);
+  uint32_t pointer_type_id = object_inst->type_id();
+  Instruction* pointer_type_inst = get_def_use_mgr()->GetDef(pointer_type_id);
+  uint32_t type_id = pointer_type_inst->GetSingleWordInOperand(1);
+  used_members_[type_id].insert(inst->GetSingleWordInOperand(1));
+}
+
+bool EliminateDeadMembersPass::RemoveDeadMembers() {
+  bool modified = false;
+
+  // First update all of the OpTypeStruct instructions.
+  get_module()->ForEachInst([&modified, this](Instruction* inst) {
+    switch (inst->opcode()) {
+      case SpvOpTypeStruct:
+        modified |= UpdateOpTypeStruct(inst);
+        break;
+      default:
+        break;
+    }
+  });
+
+  // Now update all of the instructions that reference the OpTypeStructs.
+  get_module()->ForEachInst([&modified, this](Instruction* inst) {
+    switch (inst->opcode()) {
+      case SpvOpMemberName:
+        modified |= UpdateOpMemberNameOrDecorate(inst);
+        break;
+      case SpvOpMemberDecorate:
+        modified |= UpdateOpMemberNameOrDecorate(inst);
+        break;
+      case SpvOpGroupMemberDecorate:
+        modified |= UpdateOpGroupMemberDecorate(inst);
+        break;
+      case SpvOpSpecConstantComposite:
+      case SpvOpConstantComposite:
+      case SpvOpCompositeConstruct:
+        modified |= UpdateConstantComposite(inst);
+        break;
+      case SpvOpAccessChain:
+      case SpvOpInBoundsAccessChain:
+      case SpvOpPtrAccessChain:
+      case SpvOpInBoundsPtrAccessChain:
+        modified |= UpdateAccessChain(inst);
+        break;
+      case SpvOpCompositeExtract:
+        modified |= UpdateCompsiteExtract(inst);
+        break;
+      case SpvOpCompositeInsert:
+        modified |= UpdateCompositeInsert(inst);
+        break;
+      case SpvOpArrayLength:
+        modified |= UpdateOpArrayLength(inst);
+        break;
+      case SpvOpSpecConstantOp:
+        assert(false && "Not yet implemented.");
+        // with OpCompositeExtract, OpCompositeInsert
+        // For kernels: OpAccessChain, OpInBoundsAccessChain, OpPtrAccessChain,
+        // OpInBoundsPtrAccessChain
+        break;
+      default:
+        break;
+    }
+  });
+  return modified;
+}
+
+bool EliminateDeadMembersPass::UpdateOpTypeStruct(Instruction* inst) {
+  assert(inst->opcode() == SpvOpTypeStruct);
+
+  const auto& live_members = used_members_[inst->result_id()];
+  if (live_members.size() == inst->NumInOperands()) {
+    return false;
+  }
+
+  Instruction::OperandList new_operands;
+  for (uint32_t idx : live_members) {
+    new_operands.emplace_back(inst->GetInOperand(idx));
+  }
+
+  inst->SetInOperands(std::move(new_operands));
+  context()->UpdateDefUse(inst);
+  return true;
+}
+
+bool EliminateDeadMembersPass::UpdateOpMemberNameOrDecorate(Instruction* inst) {
+  assert(inst->opcode() == SpvOpMemberName ||
+         inst->opcode() == SpvOpMemberDecorate);
+
+  uint32_t type_id = inst->GetSingleWordInOperand(0);
+  auto live_members = used_members_.find(type_id);
+  if (live_members == used_members_.end()) {
+    return false;
+  }
+
+  uint32_t orig_member_idx = inst->GetSingleWordInOperand(1);
+  uint32_t new_member_idx = GetNewMemberIndex(type_id, orig_member_idx);
+
+  if (new_member_idx == kRemovedMember) {
+    context()->KillInst(inst);
+    return true;
+  }
+
+  if (new_member_idx == orig_member_idx) {
+    return false;
+  }
+
+  inst->SetInOperand(1, {new_member_idx});
+  return true;
+}
+
+bool EliminateDeadMembersPass::UpdateOpGroupMemberDecorate(Instruction* inst) {
+  assert(inst->opcode() == SpvOpGroupMemberDecorate);
+
+  bool modified = false;
+
+  Instruction::OperandList new_operands;
+  new_operands.emplace_back(inst->GetInOperand(0));
+  for (uint32_t i = 1; i < inst->NumInOperands(); i += 2) {
+    uint32_t type_id = inst->GetSingleWordInOperand(i);
+    uint32_t member_idx = inst->GetSingleWordInOperand(i + 1);
+    uint32_t new_member_idx = GetNewMemberIndex(type_id, member_idx);
+
+    if (new_member_idx == kRemovedMember) {
+      modified = true;
+      continue;
+    }
+
+    new_operands.emplace_back(inst->GetOperand(i));
+    if (new_member_idx != member_idx) {
+      new_operands.emplace_back(
+          Operand({SPV_OPERAND_TYPE_LITERAL_INTEGER, {new_member_idx}}));
+      modified = true;
+    } else {
+      new_operands.emplace_back(inst->GetOperand(i + 1));
+    }
+  }
+
+  if (!modified) {
+    return false;
+  }
+
+  if (new_operands.size() == 1) {
+    context()->KillInst(inst);
+    return true;
+  }
+
+  inst->SetInOperands(std::move(new_operands));
+  context()->UpdateDefUse(inst);
+  return true;
+}
+
+bool EliminateDeadMembersPass::UpdateConstantComposite(Instruction* inst) {
+  assert(inst->opcode() == SpvOpConstantComposite ||
+         inst->opcode() == SpvOpCompositeConstruct);
+  uint32_t type_id = inst->type_id();
+
+  bool modified = false;
+  Instruction::OperandList new_operands;
+  for (uint32_t i = 0; i < inst->NumInOperands(); ++i) {
+    uint32_t new_idx = GetNewMemberIndex(type_id, i);
+    if (new_idx == kRemovedMember) {
+      modified = true;
+    } else {
+      new_operands.emplace_back(inst->GetInOperand(i));
+    }
+  }
+  inst->SetInOperands(std::move(new_operands));
+  context()->UpdateDefUse(inst);
+  return modified;
+}
+
+bool EliminateDeadMembersPass::UpdateAccessChain(Instruction* inst) {
+  assert(inst->opcode() == SpvOpAccessChain ||
+         inst->opcode() == SpvOpInBoundsAccessChain ||
+         inst->opcode() == SpvOpPtrAccessChain ||
+         inst->opcode() == SpvOpInBoundsPtrAccessChain);
+
+  uint32_t pointer_id = inst->GetSingleWordInOperand(0);
+  Instruction* pointer_inst = get_def_use_mgr()->GetDef(pointer_id);
+  uint32_t pointer_type_id = pointer_inst->type_id();
+  Instruction* pointer_type_inst = get_def_use_mgr()->GetDef(pointer_type_id);
+  uint32_t type_id = pointer_type_inst->GetSingleWordInOperand(1);
+
+  analysis::ConstantManager* const_mgr = context()->get_constant_mgr();
+  Instruction::OperandList new_operands;
+  bool modified = false;
+  new_operands.emplace_back(inst->GetInOperand(0));
+
+  // For pointer access chains we want to copy the element operand.
+  if (inst->opcode() == SpvOpPtrAccessChain ||
+      inst->opcode() == SpvOpInBoundsPtrAccessChain) {
+    new_operands.emplace_back(inst->GetInOperand(1));
+  }
+
+  for (uint32_t i = static_cast<uint32_t>(new_operands.size());
+       i < inst->NumInOperands(); ++i) {
+    Instruction* type_inst = get_def_use_mgr()->GetDef(type_id);
+    switch (type_inst->opcode()) {
+      case SpvOpTypeStruct: {
+        const analysis::IntConstant* member_idx =
+            const_mgr->FindDeclaredConstant(inst->GetSingleWordInOperand(i))
+                ->AsIntConstant();
+        assert(member_idx);
+        uint32_t orig_member_idx;
+        if (member_idx->type()->AsInteger()->width() == 32) {
+          orig_member_idx = member_idx->GetU32();
+        } else {
+          orig_member_idx = static_cast<uint32_t>(member_idx->GetU64());
+        }
+        uint32_t new_member_idx = GetNewMemberIndex(type_id, orig_member_idx);
+        assert(new_member_idx != kRemovedMember);
+        if (orig_member_idx != new_member_idx) {
+          InstructionBuilder ir_builder(
+              context(), inst,
+              IRContext::kAnalysisDefUse |
+                  IRContext::kAnalysisInstrToBlockMapping);
+          uint32_t const_id =
+              ir_builder.GetUintConstant(new_member_idx)->result_id();
+          new_operands.emplace_back(Operand({SPV_OPERAND_TYPE_ID, {const_id}}));
+          modified = true;
+        } else {
+          new_operands.emplace_back(inst->GetInOperand(i));
+        }
+        // The type will have already been rewritten, so use the new member
+        // index.
+        type_id = type_inst->GetSingleWordInOperand(new_member_idx);
+      } break;
+      case SpvOpTypeArray:
+      case SpvOpTypeRuntimeArray:
+      case SpvOpTypeVector:
+      case SpvOpTypeMatrix:
+        new_operands.emplace_back(inst->GetInOperand(i));
+        type_id = type_inst->GetSingleWordInOperand(0);
+        break;
+      default:
+        assert(false);
+        break;
+    }
+  }
+
+  if (!modified) {
+    return false;
+  }
+  inst->SetInOperands(std::move(new_operands));
+  context()->UpdateDefUse(inst);
+  return true;
+}
+
+uint32_t EliminateDeadMembersPass::GetNewMemberIndex(uint32_t type_id,
+                                                     uint32_t member_idx) {
+  auto live_members = used_members_.find(type_id);
+  if (live_members == used_members_.end()) {
+    return member_idx;
+  }
+
+  auto current_member = live_members->second.find(member_idx);
+  if (current_member == live_members->second.end()) {
+    return kRemovedMember;
+  }
+
+  return static_cast<uint32_t>(
+      std::distance(live_members->second.begin(), current_member));
+}
+
+bool EliminateDeadMembersPass::UpdateCompsiteExtract(Instruction* inst) {
+  uint32_t object_id = inst->GetSingleWordInOperand(0);
+  Instruction* object_inst = get_def_use_mgr()->GetDef(object_id);
+  uint32_t type_id = object_inst->type_id();
+
+  Instruction::OperandList new_operands;
+  bool modified = false;
+  new_operands.emplace_back(inst->GetInOperand(0));
+  for (uint32_t i = 1; i < inst->NumInOperands(); ++i) {
+    uint32_t member_idx = inst->GetSingleWordInOperand(i);
+    uint32_t new_member_idx = GetNewMemberIndex(type_id, member_idx);
+    assert(new_member_idx != kRemovedMember);
+    if (member_idx != new_member_idx) {
+      modified = true;
+    }
+    new_operands.emplace_back(
+        Operand({SPV_OPERAND_TYPE_LITERAL_INTEGER, {new_member_idx}}));
+
+    Instruction* type_inst = get_def_use_mgr()->GetDef(type_id);
+    switch (type_inst->opcode()) {
+      case SpvOpTypeStruct:
+        assert(i != 1 || (inst->opcode() != SpvOpPtrAccessChain &&
+                          inst->opcode() != SpvOpInBoundsPtrAccessChain));
+        // The type will have already been rewriten, so use the new member
+        // index.
+        type_id = type_inst->GetSingleWordInOperand(new_member_idx);
+        break;
+      case SpvOpTypeArray:
+      case SpvOpTypeRuntimeArray:
+      case SpvOpTypeVector:
+      case SpvOpTypeMatrix:
+        type_id = type_inst->GetSingleWordInOperand(0);
+        break;
+      default:
+        assert(false);
+    }
+  }
+
+  if (!modified) {
+    return false;
+  }
+  inst->SetInOperands(std::move(new_operands));
+  context()->UpdateDefUse(inst);
+  return true;
+}
+
+bool EliminateDeadMembersPass::UpdateCompositeInsert(Instruction* inst) {
+  uint32_t composite_id = inst->GetSingleWordInOperand(1);
+  Instruction* composite_inst = get_def_use_mgr()->GetDef(composite_id);
+  uint32_t type_id = composite_inst->type_id();
+
+  Instruction::OperandList new_operands;
+  bool modified = false;
+  new_operands.emplace_back(inst->GetInOperand(0));
+  new_operands.emplace_back(inst->GetInOperand(1));
+  for (uint32_t i = 2; i < inst->NumInOperands(); ++i) {
+    uint32_t member_idx = inst->GetSingleWordInOperand(i);
+    uint32_t new_member_idx = GetNewMemberIndex(type_id, member_idx);
+    if (new_member_idx == kRemovedMember) {
+      context()->KillInst(inst);
+      return true;
+    }
+
+    if (member_idx != new_member_idx) {
+      modified = true;
+    }
+    new_operands.emplace_back(
+        Operand({SPV_OPERAND_TYPE_LITERAL_INTEGER, {new_member_idx}}));
+
+    Instruction* type_inst = get_def_use_mgr()->GetDef(type_id);
+    switch (type_inst->opcode()) {
+      case SpvOpTypeStruct:
+        // The type will have already been rewritten, so use the new member
+        // index.
+        type_id = type_inst->GetSingleWordInOperand(new_member_idx);
+        break;
+      case SpvOpTypeArray:
+      case SpvOpTypeRuntimeArray:
+      case SpvOpTypeVector:
+      case SpvOpTypeMatrix:
+        type_id = type_inst->GetSingleWordInOperand(0);
+        break;
+      default:
+        assert(false);
+    }
+  }
+
+  if (!modified) {
+    return false;
+  }
+  inst->SetInOperands(std::move(new_operands));
+  context()->UpdateDefUse(inst);
+  return true;
+}
+
+bool EliminateDeadMembersPass::UpdateOpArrayLength(Instruction* inst) {
+  uint32_t struct_id = inst->GetSingleWordInOperand(0);
+  Instruction* struct_inst = get_def_use_mgr()->GetDef(struct_id);
+  uint32_t pointer_type_id = struct_inst->type_id();
+  Instruction* pointer_type_inst = get_def_use_mgr()->GetDef(pointer_type_id);
+  uint32_t type_id = pointer_type_inst->GetSingleWordInOperand(1);
+
+  uint32_t member_idx = inst->GetSingleWordInOperand(1);
+  uint32_t new_member_idx = GetNewMemberIndex(type_id, member_idx);
+  assert(new_member_idx != kRemovedMember);
+
+  if (member_idx == new_member_idx) {
+    return false;
+  }
+
+  inst->SetInOperand(1, {new_member_idx});
+  context()->UpdateDefUse(inst);
+  return true;
+}
+
+void EliminateDeadMembersPass::MarkStructOperandsAsFullyUsed(
+    const Instruction* inst) {
+  if (inst->type_id() != 0) {
+    MarkTypeAsFullyUsed(inst->type_id());
+  }
+
+  inst->ForEachInId([this](const uint32_t* id) {
+    Instruction* instruction = get_def_use_mgr()->GetDef(*id);
+    if (instruction->type_id() != 0) {
+      MarkTypeAsFullyUsed(instruction->type_id());
+    }
+  });
+}
+
+}  // namespace opt
+}  // namespace spvtools

3rdparty/spirv-tools/source/opt/eliminate_dead_members_pass.h

@@ -0,0 +1,145 @@
+// Copyright (c) 2019 Google LLC
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+#ifndef SOURCE_OPT_ELIMINATE_DEAD_MEMBERS_PASS_H_
+#define SOURCE_OPT_ELIMINATE_DEAD_MEMBERS_PASS_H_
+
+#include "source/opt/def_use_manager.h"
+#include "source/opt/function.h"
+#include "source/opt/mem_pass.h"
+#include "source/opt/module.h"
+
+namespace spvtools {
+namespace opt {
+
+// Remove unused members from structures.  The remaining members will remain at
+// the same offset.
+class EliminateDeadMembersPass : public MemPass {
+ public:
+  const char* name() const override { return "eliminate-dead-members"; }
+  Status Process() override;
+
+  IRContext::Analysis GetPreservedAnalyses() override {
+    return IRContext::kAnalysisDefUse |
+           IRContext::kAnalysisInstrToBlockMapping |
+           IRContext::kAnalysisCombinators | IRContext::kAnalysisCFG |
+           IRContext::kAnalysisDominatorAnalysis |
+           IRContext::kAnalysisLoopAnalysis |
+           IRContext::kAnalysisScalarEvolution |
+           IRContext::kAnalysisRegisterPressure |
+           IRContext::kAnalysisValueNumberTable |
+           IRContext::kAnalysisStructuredCFG |
+           IRContext::kAnalysisBuiltinVarId |
+           IRContext::kAnalysisIdToFuncMapping;
+  }
+
+ private:
+  // Populate |used_members_| with the member of structures that are live in the
+  // current context.
+  void FindLiveMembers();
+
+  // Add to |used_members_| the member of structures that are live in
+  // |function|.
+  void FindLiveMembers(const Function& function);
+  // Add to |used_members_| the member of structures that are live in |inst|.
+  void FindLiveMembers(const Instruction* inst);
+
+  // Add to |used_members_| the members that are live in the |OpStore|
+  // instruction |inst|.
+  void MarkMembersAsLiveForStore(const Instruction* inst);
+
+  // Add to |used_members_| the members that are live in the |OpCopyMemory*|
+  // instruction |inst|.
+  void MarkMembersAsLiveForCopyMemory(const Instruction* inst);
+
+  // Add to |used_members_| the members that are live in the
+  // |OpCompositeExtract| instruction |inst|.
+  void MarkMembersAsLiveForExtract(const Instruction* inst);
+
+  // Add to |used_members_| the members that are live in the |Op*AccessChain|
+  // instruction |inst|.
+  void MarkMembersAsLiveForAccessChain(const Instruction* inst);
+
+  // Add the member referenced by the OpArrayLength instruction |inst| to
+  // |uses_members_|.
+  void MarkMembersAsLiveForArrayLength(const Instruction* inst);
+
+  // Remove dead members from structs and updates any instructions that need to
+  // be updated as a consequence.  Return true if something changed.
+  bool RemoveDeadMembers();
+
+  // Update |inst|, which must be an |OpMemberName| or |OpMemberDecorate|
+  // instruction, so it references the correct member after the struct is
+  // updated.  Return true if something changed.
+  bool UpdateOpMemberNameOrDecorate(Instruction* inst);
+
+  // Update |inst|, which must be an |OpGroupMemberDecorate| instruction, so it
+  // references the correct member after the struct is updated.  Return true if
+  // something changed.
+  bool UpdateOpGroupMemberDecorate(Instruction* inst);
+
+  // Update the |OpTypeStruct| instruction |inst| my removing the members that
+  // are not live.  Return true if something changed.
+  bool UpdateOpTypeStruct(Instruction* inst);
+
+  // Update the |OpConstantComposite| instruction |inst| to match the change
+  // made to the type that was being generated.  Return true if something
+  // changed.
+  bool UpdateConstantComposite(Instruction* inst);
+
+  // Update the |Op*AccessChain| instruction |inst| to reference the correct
+  // members. All members referenced in the access chain must be live.  This
+  // function must be called after the |OpTypeStruct| instruction for the type
+  // has been updated.  Return true if something changed.
+  bool UpdateAccessChain(Instruction* inst);
+
+  // Update the |OpCompositeExtract| instruction |inst| to reference the correct
+  // members. All members referenced in the instruction must be live.  This
+  // function must be called after the |OpTypeStruct| instruction for the type
+  // has been updated.  Return true if something changed.
+  bool UpdateCompsiteExtract(Instruction* inst);
+
+  // Update the |OpCompositeInsert| instruction |inst| to reference the correct
+  // members. If the member being inserted is not live, then |inst| is killed.
+  // This function must be called after the |OpTypeStruct| instruction for the
+  // type has been updated.  Return true if something changed.
+  bool UpdateCompositeInsert(Instruction* inst);
+
+  // Update the |OpArrayLength| instruction |inst| to reference the correct
+  // member. The member referenced in the instruction must be live.  Return true
+  // if something changed.
+  bool UpdateOpArrayLength(Instruction* inst);
+
+  // Add all of the members of type |type_id| and members of any subtypes to
+  // |used_members_|.
+  void MarkTypeAsFullyUsed(uint32_t type_id);
+
+  // Add all of the members of the type of the operand |in_idx| in |inst| and
+  // members of any subtypes to |uses_members_|.
+  void MarkOperandTypeAsFullyUsed(const Instruction* inst, uint32_t in_idx);
+
+  // Return the index of the member that use to be the |member_idx|th member of
+  // |type_id|.  If the member has been removed, |kRemovedMember| is returned.
+  uint32_t GetNewMemberIndex(uint32_t type_id, uint32_t member_idx);
+
+  // A map from a type id to a set of indices representing the members of the
+  // type that are used, and must be kept.
+  std::unordered_map<uint32_t, std::set<uint32_t>> used_members_;
+  void MarkStructOperandsAsFullyUsed(const Instruction* inst);
+};
+
+}  // namespace opt
+}  // namespace spvtools
+
+#endif  // SOURCE_OPT_ELIMINATE_DEAD_MEMBERS_PASS_H_

3rdparty/spirv-tools/source/opt/instrument_pass.cpp

@@ -614,8 +614,14 @@ bool InstrumentPass::InstrumentFunction(Function* func, uint32_t stage_idx,
     ++function_idx;
   }
   std::vector<std::unique_ptr<BasicBlock>> new_blks;
-  // Start count after function instruction
+  // Start count after function and param instructions
   uint32_t instruction_idx = funcIdx2offset_[function_idx] + 1;
+  func->ForEachParam(
+      [this, &instruction_idx](const Instruction* i) {
+        (void)i;
+        ++instruction_idx;
+      },
+      true);
   // Using block iterators here because of block erasures and insertions.
   for (auto bi = func->begin(); bi != func->end(); ++bi) {
     // Count block's label
@@ -784,8 +790,14 @@ void InstrumentPass::InitializeInstrument() {
   auto prev_fn = get_module()->begin();
   auto curr_fn = prev_fn;
   for (++curr_fn; curr_fn != get_module()->end(); ++curr_fn) {
-    // Count function and end instructions
+    // Count function, end and param instructions
     uint32_t func_size = 2;
+    prev_fn->ForEachParam(
+        [this, &func_size](const Instruction* i) {
+          (void)i;
+          ++func_size;
+        },
+        true);
     for (auto& blk : *prev_fn) {
       // Count label
       func_size += 1;
@@ -794,7 +806,7 @@ void InstrumentPass::InitializeInstrument() {
         func_size += static_cast<uint32_t>(inst.dbg_line_insts().size());
       }
     }
-    funcIdx2offset_[func_idx] = func_size;
+    funcIdx2offset_[func_idx] = funcIdx2offset_[func_idx - 1] + func_size;
     ++prev_fn;
     ++func_idx;
   }

3rdparty/spirv-tools/source/opt/optimizer.cpp

@@ -332,6 +332,8 @@ bool Optimizer::RegisterPassFromFlag(const std::string& flag) {
     RegisterPass(CreateDeadInsertElimPass());
   } else if (pass_name == "eliminate-dead-variables") {
     RegisterPass(CreateDeadVariableEliminationPass());
+  } else if (pass_name == "eliminate-dead-members") {
+    RegisterPass(CreateEliminateDeadMembersPass());
   } else if (pass_name == "fold-spec-const-op-composite") {
     RegisterPass(CreateFoldSpecConstantOpAndCompositePass());
   } else if (pass_name == "loop-unswitch") {
@@ -542,6 +544,11 @@ Optimizer::PassToken CreateEliminateDeadFunctionsPass() {
       MakeUnique<opt::EliminateDeadFunctionsPass>());
 }
 
+Optimizer::PassToken CreateEliminateDeadMembersPass() {
+  return MakeUnique<Optimizer::PassToken::Impl>(
+      MakeUnique<opt::EliminateDeadMembersPass>());
+}
+
 Optimizer::PassToken CreateSetSpecConstantDefaultValuePass(
     const std::unordered_map<uint32_t, std::string>& id_value_map) {
   return MakeUnique<Optimizer::PassToken::Impl>(

3rdparty/spirv-tools/source/opt/passes.h

@@ -31,6 +31,7 @@
 #include "source/opt/dead_variable_elimination.h"
 #include "source/opt/eliminate_dead_constant_pass.h"
 #include "source/opt/eliminate_dead_functions_pass.h"
+#include "source/opt/eliminate_dead_members_pass.h"
 #include "source/opt/flatten_decoration_pass.h"
 #include "source/opt/fold_spec_constant_op_and_composite_pass.h"
 #include "source/opt/freeze_spec_constant_value_pass.h"

3rdparty/spirv-tools/source/reduce/merge_blocks_reduction_opportunity.cpp

@@ -34,9 +34,21 @@ MergeBlocksReductionOpportunity::MergeBlocksReductionOpportunity(
 }
 
 bool MergeBlocksReductionOpportunity::PreconditionHolds() {
-  // By construction, it is not possible for the merging of A->B to disable the
-  // merging of C->D, even when B and C are the same block.
-  return true;
+  // Merge block opportunities can disable each other.
+  // Example: Given blocks: A->B->C.
+  // A is a loop header; B and C are blocks in the loop; C ends with OpReturn.
+  // There are two opportunities: B and C can be merged with their predecessors.
+  // Merge C. B now ends with OpReturn. We now just have: A->B.
+  // Merge B is now disabled, as this would lead to A, a loop header, ending
+  // with an OpReturn, which is invalid.
+
+  const auto predecessors = context_->cfg()->preds(successor_block_->id());
+  assert(1 == predecessors.size() &&
+         "For a successor to be merged into its predecessor, exactly one "
+         "predecessor must be present.");
+  const uint32_t predecessor_id = predecessors[0];
+  BasicBlock* predecessor_block = context_->get_instr_block(predecessor_id);
+  return blockmergeutil::CanMergeWithSuccessor(context_, predecessor_block);
 }
 
 void MergeBlocksReductionOpportunity::Apply() {
@@ -50,6 +62,7 @@ void MergeBlocksReductionOpportunity::Apply() {
          "predecessor must be present.");
   const uint32_t predecessor_id = predecessors[0];
 
+  // We need an iterator pointing to the predecessor, hence the loop.
   for (auto bi = function_->begin(); bi != function_->end(); ++bi) {
     if (bi->id() == predecessor_id) {
       blockmergeutil::MergeWithSuccessor(context_, function_, bi);

3rdparty/spirv-tools/source/val/validate.cpp

@@ -385,7 +385,6 @@ spv_result_t ValidateBinaryUsingContextAndValidationState(
       Instruction* inst = const_cast<Instruction*>(&instruction);
       vstate->RegisterInstruction(inst);
     }
-    if (auto error = UpdateIdUse(*vstate, &instruction)) return error;
   }
 
   if (!vstate->has_memory_model_specified())
@@ -399,6 +398,19 @@ spv_result_t ValidateBinaryUsingContextAndValidationState(
   // Catch undefined forward references before performing further checks.
   if (auto error = ValidateForwardDecls(*vstate)) return error;
 
+  // ID usage needs be handled in its own iteration of the instructions,
+  // between the two others. It depends on the first loop to have been
+  // finished, so that all instructions have been registered. And the following
+  // loop depends on all of the usage data being populated. Thus it cannot live
+  // in either of those iterations.
+  // It should also live after the forward declaration check, since it will
+  // have problems with missing forward declarations, but give less useful error
+  // messages.
+  for (size_t i = 0; i < vstate->ordered_instructions().size(); ++i) {
+    auto& instruction = vstate->ordered_instructions()[i];
+    if (auto error = UpdateIdUse(*vstate, &instruction)) return error;
+  }
+
   // Validate individual opcodes.
   for (size_t i = 0; i < vstate->ordered_instructions().size(); ++i) {
     auto& instruction = vstate->ordered_instructions()[i];

3rdparty/spirv-tools/source/val/validate_decorations.cpp

@@ -1185,6 +1185,9 @@ spv_result_t CheckFPRoundingModeForShaders(ValidationState_t& vstate,
   // Validates Object operand of an OpStore
   for (const auto& use : inst.uses()) {
     const auto store = use.first;
+    if (store->opcode() == SpvOpFConvert) continue;
+    if (spvOpcodeIsDebug(store->opcode())) continue;
+    if (spvOpcodeIsDecoration(store->opcode())) continue;
     if (store->opcode() != SpvOpStore) {
       return vstate.diag(SPV_ERROR_INVALID_ID, &inst)
              << "FPRoundingMode decoration can be applied only to the "

3rdparty/spirv-tools/source/val/validate_function.cpp

@@ -41,18 +41,22 @@ spv_result_t ValidateFunction(ValidationState_t& _, const Instruction* inst) {
            << _.getIdName(return_id) << "'.";
   }
 
+  const std::vector<SpvOp> acceptable = {
+      SpvOpDecorate,
+      SpvOpEnqueueKernel,
+      SpvOpEntryPoint,
+      SpvOpExecutionMode,
+      SpvOpExecutionModeId,
+      SpvOpFunctionCall,
+      SpvOpGetKernelNDrangeSubGroupCount,
+      SpvOpGetKernelNDrangeMaxSubGroupSize,
+      SpvOpGetKernelWorkGroupSize,
+      SpvOpGetKernelPreferredWorkGroupSizeMultiple,
+      SpvOpGetKernelLocalSizeForSubgroupCount,
+      SpvOpGetKernelMaxNumSubgroups,
+      SpvOpName};
   for (auto& pair : inst->uses()) {
     const auto* use = pair.first;
-    const std::vector<SpvOp> acceptable = {
-        SpvOpFunctionCall,
-        SpvOpEntryPoint,
-        SpvOpEnqueueKernel,
-        SpvOpGetKernelNDrangeSubGroupCount,
-        SpvOpGetKernelNDrangeMaxSubGroupSize,
-        SpvOpGetKernelWorkGroupSize,
-        SpvOpGetKernelPreferredWorkGroupSizeMultiple,
-        SpvOpGetKernelLocalSizeForSubgroupCount,
-        SpvOpGetKernelMaxNumSubgroups};
     if (std::find(acceptable.begin(), acceptable.end(), use->opcode()) ==
         acceptable.end()) {
       return _.diag(SPV_ERROR_INVALID_ID, use)

3rdparty/spirv-tools/source/val/validate_type.cpp

@@ -324,10 +324,12 @@ spv_result_t ValidateTypeFunction(ValidationState_t& _,
            << num_args << " arguments.";
   }
 
-  // The only valid uses of OpTypeFunction are in an OpFunction instruction.
+  // The only valid uses of OpTypeFunction are in an OpFunction, debugging, or
+  // decoration instruction.
   for (auto& pair : inst->uses()) {
     const auto* use = pair.first;
-    if (use->opcode() != SpvOpFunction) {
+    if (use->opcode() != SpvOpFunction && !spvOpcodeIsDebug(use->opcode()) &&
+        !spvOpcodeIsDecoration(use->opcode())) {
       return _.diag(SPV_ERROR_INVALID_ID, use)
              << "Invalid use of function type result id "
              << _.getIdName(inst->id()) << ".";

3rdparty/spirv-tools/test/opt/CMakeLists.txt

@@ -34,6 +34,7 @@ add_spvtools_unittest(TARGET opt
        def_use_test.cpp
        eliminate_dead_const_test.cpp
        eliminate_dead_functions_test.cpp
+       eliminate_dead_member_test.cpp
        feature_manager_test.cpp
        flatten_decoration_test.cpp
        fold_spec_const_op_composite_test.cpp

3rdparty/spirv-tools/test/opt/block_merge_test.cpp

@@ -483,6 +483,8 @@ TEST_F(BlockMergeTest, RemoveStructuredDeclaration) {
 ; CHECK-NOT: OpLoopMerge
 ; CHECK: OpReturn
 ; CHECK: [[continue:%\w+]] = OpLabel
+; CHECK-NEXT: OpBranch [[block:%\w+]]
+; CHECK: [[block]] = OpLabel
 ; CHECK-NEXT: OpBranch [[header]]
 OpCapability Shader
 %1 = OpExtInstImport "GLSL.std.450"
@@ -880,6 +882,47 @@ OpFunctionEnd
                                         prefix + suffix_after, true, true);
 }
 
+TEST_F(BlockMergeTest, UnreachableLoop) {
+  const std::string spirv = R"(OpCapability Shader
+%1 = OpExtInstImport "GLSL.std.450"
+OpMemoryModel Logical GLSL450
+OpEntryPoint Fragment %main "main"
+OpExecutionMode %main OriginUpperLeft
+OpSource ESSL 310
+OpName %main "main"
+%void = OpTypeVoid
+%4 = OpTypeFunction %void
+%int = OpTypeInt 32 1
+%_ptr_Function_int = OpTypePointer Function %int
+%bool = OpTypeBool
+%false = OpConstantFalse %bool
+%main = OpFunction %void None %4
+%9 = OpLabel
+OpBranch %10
+%11 = OpLabel
+OpLoopMerge %12 %13 None
+OpBranchConditional %false %13 %14
+%13 = OpLabel
+OpSelectionMerge %15 None
+OpBranchConditional %false %16 %17
+%16 = OpLabel
+OpBranch %15
+%17 = OpLabel
+OpBranch %15
+%15 = OpLabel
+OpBranch %11
+%14 = OpLabel
+OpReturn
+%12 = OpLabel
+OpBranch %10
+%10 = OpLabel
+OpReturn
+OpFunctionEnd
+)";
+
+  SinglePassRunAndCheck<BlockMergePass>(spirv, spirv, true, true);
+}
+
 // TODO(greg-lunarg): Add tests to verify handling of these cases:
 //
 //    More complex control flow

3rdparty/spirv-tools/test/opt/eliminate_dead_member_test.cpp

@@ -0,0 +1,1031 @@
+// Copyright (c) 2019 Google LLC
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+#include "assembly_builder.h"
+#include "gmock/gmock.h"
+#include "pass_fixture.h"
+#include "pass_utils.h"
+
+namespace {
+
+using namespace spvtools;
+
+using EliminateDeadMemberTest = opt::PassTest<::testing::Test>;
+
+TEST_F(EliminateDeadMemberTest, RemoveMember1) {
+  // Test that the member "y" is removed.
+  // Update OpMemberName for |y| and |z|.
+  // Update OpMemberDecorate for |y| and |z|.
+  // Update OpAccessChain for access to |z|.
+  const std::string text = R"(
+; CHECK: OpName
+; CHECK-NEXT: OpMemberName %type__Globals 0 "x"
+; CHECK-NEXT: OpMemberName %type__Globals 1 "z"
+; CHECK-NOT: OpMemberName
+; CHECK: OpMemberDecorate %type__Globals 0 Offset 0
+; CHECK: OpMemberDecorate %type__Globals 1 Offset 8
+; CHECK: %type__Globals = OpTypeStruct %float %float
+; CHECK: OpAccessChain %_ptr_Uniform_float %_Globals %int_0
+; CHECK: OpAccessChain %_ptr_Uniform_float %_Globals %uint_1
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Vertex %main "main" %in_var_Position %gl_Position
+               OpSource HLSL 600
+               OpName %type__Globals "type.$Globals"
+               OpMemberName %type__Globals 0 "x"
+               OpMemberName %type__Globals 1 "y"
+               OpMemberName %type__Globals 2 "z"
+               OpName %_Globals "$Globals"
+               OpName %in_var_Position "in.var.Position"
+               OpName %main "main"
+               OpDecorate %gl_Position BuiltIn Position
+               OpDecorate %in_var_Position Location 0
+               OpDecorate %_Globals DescriptorSet 0
+               OpDecorate %_Globals Binding 0
+               OpMemberDecorate %type__Globals 0 Offset 0
+               OpMemberDecorate %type__Globals 1 Offset 4
+               OpMemberDecorate %type__Globals 2 Offset 8
+               OpDecorate %type__Globals Block
+        %int = OpTypeInt 32 1
+      %int_0 = OpConstant %int 0
+      %float = OpTypeFloat 32
+      %int_2 = OpConstant %int 2
+%type__Globals = OpTypeStruct %float %float %float
+%_ptr_Uniform_type__Globals = OpTypePointer Uniform %type__Globals
+    %v4float = OpTypeVector %float 4
+%_ptr_Input_v4float = OpTypePointer Input %v4float
+%_ptr_Output_v4float = OpTypePointer Output %v4float
+       %void = OpTypeVoid
+         %15 = OpTypeFunction %void
+%_ptr_Uniform_float = OpTypePointer Uniform %float
+   %_Globals = OpVariable %_ptr_Uniform_type__Globals Uniform
+%in_var_Position = OpVariable %_ptr_Input_v4float Input
+%gl_Position = OpVariable %_ptr_Output_v4float Output
+       %main = OpFunction %void None %15
+         %17 = OpLabel
+         %18 = OpLoad %v4float %in_var_Position
+         %19 = OpAccessChain %_ptr_Uniform_float %_Globals %int_0
+         %20 = OpLoad %float %19
+         %21 = OpCompositeExtract %float %18 0
+         %22 = OpFAdd %float %21 %20
+         %23 = OpCompositeInsert %v4float %22 %18 0
+         %24 = OpCompositeExtract %float %18 1
+         %25 = OpCompositeInsert %v4float %24 %23 1
+         %26 = OpAccessChain %_ptr_Uniform_float %_Globals %int_2
+         %27 = OpLoad %float %26
+         %28 = OpCompositeExtract %float %18 2
+         %29 = OpFAdd %float %28 %27
+         %30 = OpCompositeInsert %v4float %29 %25 2
+               OpStore %gl_Position %30
+               OpReturn
+               OpFunctionEnd
+)";
+
+  SinglePassRunAndMatch<opt::EliminateDeadMembersPass>(text, true);
+}
+
+TEST_F(EliminateDeadMemberTest, RemoveMemberWithGroupDecorations) {
+  // Test that the member "y" is removed.
+  // Update OpGroupMemberDecorate for %type__Globals member 1 and 2.
+  // Update OpAccessChain for access to %type__Globals member 2.
+  const std::string text = R"(
+; CHECK: OpDecorate [[gr1:%\w+]] Offset 0
+; CHECK: OpDecorate [[gr2:%\w+]] Offset 4
+; CHECK: OpDecorate [[gr3:%\w+]] Offset 8
+; CHECK: [[gr1]] = OpDecorationGroup
+; CHECK: [[gr2]] = OpDecorationGroup
+; CHECK: [[gr3]] = OpDecorationGroup
+; CHECK: OpGroupMemberDecorate [[gr1]] %type__Globals 0
+; CHECK-NOT: OpGroupMemberDecorate [[gr2]]
+; CHECK: OpGroupMemberDecorate [[gr3]] %type__Globals 1
+; CHECK: %type__Globals = OpTypeStruct %float %float
+; CHECK: OpAccessChain %_ptr_Uniform_float %_Globals %int_0
+; CHECK: OpAccessChain %_ptr_Uniform_float %_Globals %uint_1
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Vertex %main "main" %in_var_Position %gl_Position
+               OpSource HLSL 600
+               OpName %type__Globals "type.$Globals"
+               OpName %_Globals "$Globals"
+               OpDecorate %gl_Position BuiltIn Position
+               OpDecorate %in_var_Position Location 0
+               OpDecorate %_Globals DescriptorSet 0
+               OpDecorate %_Globals Binding 0
+               OpDecorate %gr1 Offset 0
+               OpDecorate %gr2 Offset 4
+               OpDecorate %gr3 Offset 8
+               OpDecorate %type__Globals Block
+        %gr1 = OpDecorationGroup
+        %gr2 = OpDecorationGroup
+        %gr3 = OpDecorationGroup
+               OpGroupMemberDecorate %gr1 %type__Globals 0
+               OpGroupMemberDecorate %gr2 %type__Globals 1
+               OpGroupMemberDecorate %gr3 %type__Globals 2
+        %int = OpTypeInt 32 1
+      %int_0 = OpConstant %int 0
+      %float = OpTypeFloat 32
+      %int_2 = OpConstant %int 2
+%type__Globals = OpTypeStruct %float %float %float
+%_ptr_Uniform_type__Globals = OpTypePointer Uniform %type__Globals
+    %v4float = OpTypeVector %float 4
+%_ptr_Input_v4float = OpTypePointer Input %v4float
+%_ptr_Output_v4float = OpTypePointer Output %v4float
+       %void = OpTypeVoid
+         %15 = OpTypeFunction %void
+%_ptr_Uniform_float = OpTypePointer Uniform %float
+   %_Globals = OpVariable %_ptr_Uniform_type__Globals Uniform
+%in_var_Position = OpVariable %_ptr_Input_v4float Input
+%gl_Position = OpVariable %_ptr_Output_v4float Output
+       %main = OpFunction %void None %15
+         %17 = OpLabel
+         %18 = OpLoad %v4float %in_var_Position
+         %19 = OpAccessChain %_ptr_Uniform_float %_Globals %int_0
+         %20 = OpLoad %float %19
+         %21 = OpCompositeExtract %float %18 0
+         %22 = OpFAdd %float %21 %20
+         %23 = OpCompositeInsert %v4float %22 %18 0
+         %24 = OpCompositeExtract %float %18 1
+         %25 = OpCompositeInsert %v4float %24 %23 1
+         %26 = OpAccessChain %_ptr_Uniform_float %_Globals %int_2
+         %27 = OpLoad %float %26
+         %28 = OpCompositeExtract %float %18 2
+         %29 = OpFAdd %float %28 %27
+         %30 = OpCompositeInsert %v4float %29 %25 2
+               OpStore %gl_Position %30
+               OpReturn
+               OpFunctionEnd
+)";
+
+  // Skipping validation because of a bug in the validator.  See issue #2376.
+  SinglePassRunAndMatch<opt::EliminateDeadMembersPass>(text, false);
+}
+
+TEST_F(EliminateDeadMemberTest, RemoveMemberUpdateConstant) {
+  // Test that the member "x" is removed.
+  // Update the OpConstantComposite instruction.
+  const std::string text = R"(
+; CHECK: OpName
+; CHECK-NEXT: OpMemberName %type__Globals 0 "y"
+; CHECK-NEXT: OpMemberName %type__Globals 1 "z"
+; CHECK-NOT: OpMemberName
+; CHECK: OpMemberDecorate %type__Globals 0 Offset 4
+; CHECK: OpMemberDecorate %type__Globals 1 Offset 8
+; CHECK: %type__Globals = OpTypeStruct %float %float
+; CHECK: OpConstantComposite %type__Globals %float_1 %float_2
+; CHECK: OpAccessChain %_ptr_Uniform_float %_Globals %uint_0
+; CHECK: OpAccessChain %_ptr_Uniform_float %_Globals %uint_1
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Vertex %main "main" %in_var_Position %gl_Position
+               OpSource HLSL 600
+               OpName %type__Globals "type.$Globals"
+               OpMemberName %type__Globals 0 "x"
+               OpMemberName %type__Globals 1 "y"
+               OpMemberName %type__Globals 2 "z"
+               OpName %_Globals "$Globals"
+               OpName %in_var_Position "in.var.Position"
+               OpName %main "main"
+               OpDecorate %gl_Position BuiltIn Position
+               OpDecorate %in_var_Position Location 0
+               OpDecorate %_Globals DescriptorSet 0
+               OpDecorate %_Globals Binding 0
+               OpMemberDecorate %type__Globals 0 Offset 0
+               OpMemberDecorate %type__Globals 1 Offset 4
+               OpMemberDecorate %type__Globals 2 Offset 8
+               OpDecorate %type__Globals Block
+        %int = OpTypeInt 32 1
+      %int_1 = OpConstant %int 1
+      %float = OpTypeFloat 32
+    %float_0 = OpConstant %float 0
+    %float_1 = OpConstant %float 1
+    %float_2 = OpConstant %float 2
+      %int_2 = OpConstant %int 2
+%type__Globals = OpTypeStruct %float %float %float
+         %13 = OpConstantComposite %type__Globals %float_0 %float_1 %float_2
+%_ptr_Uniform_type__Globals = OpTypePointer Uniform %type__Globals
+    %v4float = OpTypeVector %float 4
+%_ptr_Input_v4float = OpTypePointer Input %v4float
+%_ptr_Output_v4float = OpTypePointer Output %v4float
+       %void = OpTypeVoid
+         %19 = OpTypeFunction %void
+%_ptr_Uniform_float = OpTypePointer Uniform %float
+   %_Globals = OpVariable %_ptr_Uniform_type__Globals Uniform
+%in_var_Position = OpVariable %_ptr_Input_v4float Input
+%gl_Position = OpVariable %_ptr_Output_v4float Output
+       %main = OpFunction %void None %19
+         %21 = OpLabel
+         %22 = OpLoad %v4float %in_var_Position
+         %23 = OpAccessChain %_ptr_Uniform_float %_Globals %int_1
+         %24 = OpLoad %float %23
+         %25 = OpCompositeExtract %float %22 0
+         %26 = OpFAdd %float %25 %24
+         %27 = OpCompositeInsert %v4float %26 %22 0
+         %28 = OpCompositeExtract %float %22 1
+         %29 = OpCompositeInsert %v4float %28 %27 1
+         %30 = OpAccessChain %_ptr_Uniform_float %_Globals %int_2
+         %31 = OpLoad %float %30
+         %32 = OpCompositeExtract %float %22 2
+         %33 = OpFAdd %float %32 %31
+         %34 = OpCompositeInsert %v4float %33 %29 2
+               OpStore %gl_Position %34
+               OpReturn
+               OpFunctionEnd
+)";
+
+  SinglePassRunAndMatch<opt::EliminateDeadMembersPass>(text, true);
+}
+
+TEST_F(EliminateDeadMemberTest, RemoveMemberUpdateCompositeConstruct) {
+  // Test that the member "x" is removed.
+  // Update the OpConstantComposite instruction.
+  const std::string text = R"(
+; CHECK: OpName
+; CHECK-NEXT: OpMemberName %type__Globals 0 "y"
+; CHECK-NEXT: OpMemberName %type__Globals 1 "z"
+; CHECK-NOT: OpMemberName
+; CHECK: OpMemberDecorate %type__Globals 0 Offset 4
+; CHECK: OpMemberDecorate %type__Globals 1 Offset 8
+; CHECK: %type__Globals = OpTypeStruct %float %float
+; CHECK: OpCompositeConstruct %type__Globals %float_1 %float_2
+; CHECK: OpAccessChain %_ptr_Uniform_float %_Globals %uint_0
+; CHECK: OpAccessChain %_ptr_Uniform_float %_Globals %uint_1
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Vertex %main "main" %in_var_Position %gl_Position
+               OpSource HLSL 600
+               OpName %type__Globals "type.$Globals"
+               OpMemberName %type__Globals 0 "x"
+               OpMemberName %type__Globals 1 "y"
+               OpMemberName %type__Globals 2 "z"
+               OpName %_Globals "$Globals"
+               OpName %in_var_Position "in.var.Position"
+               OpName %main "main"
+               OpDecorate %gl_Position BuiltIn Position
+               OpDecorate %in_var_Position Location 0
+               OpDecorate %_Globals DescriptorSet 0
+               OpDecorate %_Globals Binding 0
+               OpMemberDecorate %type__Globals 0 Offset 0
+               OpMemberDecorate %type__Globals 1 Offset 4
+               OpMemberDecorate %type__Globals 2 Offset 8
+               OpDecorate %type__Globals Block
+        %int = OpTypeInt 32 1
+      %int_1 = OpConstant %int 1
+      %float = OpTypeFloat 32
+    %float_0 = OpConstant %float 0
+    %float_1 = OpConstant %float 1
+    %float_2 = OpConstant %float 2
+      %int_2 = OpConstant %int 2
+%type__Globals = OpTypeStruct %float %float %float
+%_ptr_Uniform_type__Globals = OpTypePointer Uniform %type__Globals
+    %v4float = OpTypeVector %float 4
+%_ptr_Input_v4float = OpTypePointer Input %v4float
+%_ptr_Output_v4float = OpTypePointer Output %v4float
+       %void = OpTypeVoid
+         %19 = OpTypeFunction %void
+%_ptr_Uniform_float = OpTypePointer Uniform %float
+   %_Globals = OpVariable %_ptr_Uniform_type__Globals Uniform
+%in_var_Position = OpVariable %_ptr_Input_v4float Input
+%gl_Position = OpVariable %_ptr_Output_v4float Output
+       %main = OpFunction %void None %19
+         %21 = OpLabel
+         %13 = OpCompositeConstruct %type__Globals %float_0 %float_1 %float_2
+         %22 = OpLoad %v4float %in_var_Position
+         %23 = OpAccessChain %_ptr_Uniform_float %_Globals %int_1
+         %24 = OpLoad %float %23
+         %25 = OpCompositeExtract %float %22 0
+         %26 = OpFAdd %float %25 %24
+         %27 = OpCompositeInsert %v4float %26 %22 0
+         %28 = OpCompositeExtract %float %22 1
+         %29 = OpCompositeInsert %v4float %28 %27 1
+         %30 = OpAccessChain %_ptr_Uniform_float %_Globals %int_2
+         %31 = OpLoad %float %30
+         %32 = OpCompositeExtract %float %22 2
+         %33 = OpFAdd %float %32 %31
+         %34 = OpCompositeInsert %v4float %33 %29 2
+               OpStore %gl_Position %34
+               OpReturn
+               OpFunctionEnd
+)";
+
+  SinglePassRunAndMatch<opt::EliminateDeadMembersPass>(text, true);
+}
+
+TEST_F(EliminateDeadMemberTest, RemoveMembersUpdateInserExtract1) {
+  // Test that the members "x" and "z" are removed.
+  // Update the OpCompositeExtract instruction.
+  // Remove the OpCompositeInsert instruction since the member being inserted is
+  // dead.
+  const std::string text = R"(
+; CHECK: OpName
+; CHECK-NEXT: OpMemberName %type__Globals 0 "y"
+; CHECK-NOT: OpMemberName
+; CHECK: OpMemberDecorate %type__Globals 0 Offset 4
+; CHECK-NOT: OpMemberDecorate %type__Globals 1 Offset
+; CHECK: %type__Globals = OpTypeStruct %float
+; CHECK: [[ld:%\w+]] = OpLoad %type__Globals %_Globals
+; CHECK: OpCompositeExtract %float [[ld]] 0
+; CHECK-NOT: OpCompositeInsert
+; CHECK: OpReturn
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Vertex %main "main"
+               OpSource HLSL 600
+               OpName %type__Globals "type.$Globals"
+               OpMemberName %type__Globals 0 "x"
+               OpMemberName %type__Globals 1 "y"
+               OpMemberName %type__Globals 2 "z"
+               OpName %_Globals "$Globals"
+               OpName %main "main"
+               OpDecorate %_Globals DescriptorSet 0
+               OpDecorate %_Globals Binding 0
+               OpMemberDecorate %type__Globals 0 Offset 0
+               OpMemberDecorate %type__Globals 1 Offset 4
+               OpMemberDecorate %type__Globals 2 Offset 8
+               OpDecorate %type__Globals Block
+      %float = OpTypeFloat 32
+%type__Globals = OpTypeStruct %float %float %float
+%_ptr_Uniform_type__Globals = OpTypePointer Uniform %type__Globals
+       %void = OpTypeVoid
+          %7 = OpTypeFunction %void
+   %_Globals = OpVariable %_ptr_Uniform_type__Globals Uniform
+       %main = OpFunction %void None %7
+          %8 = OpLabel
+          %9 = OpLoad %type__Globals %_Globals
+         %10 = OpCompositeExtract %float %9 1
+         %11 = OpCompositeInsert %type__Globals %10 %9 2
+               OpReturn
+               OpFunctionEnd
+
+)";
+
+  SinglePassRunAndMatch<opt::EliminateDeadMembersPass>(text, true);
+}
+
+TEST_F(EliminateDeadMemberTest, RemoveMembersUpdateInserExtract2) {
+  // Test that the members "x" and "z" are removed.
+  // Update the OpCompositeExtract instruction.
+  // Update the OpCompositeInsert instruction.
+  const std::string text = R"(
+; CHECK: OpName
+; CHECK-NEXT: OpMemberName %type__Globals 0 "y"
+; CHECK-NOT: OpMemberName
+; CHECK: OpMemberDecorate %type__Globals 0 Offset 4
+; CHECK-NOT: OpMemberDecorate %type__Globals 1 Offset
+; CHECK: %type__Globals = OpTypeStruct %float
+; CHECK: [[ld:%\w+]] = OpLoad %type__Globals %_Globals
+; CHECK: [[ex:%\w+]] = OpCompositeExtract %float [[ld]] 0
+; CHECK: OpCompositeInsert %type__Globals [[ex]] [[ld]] 0
+; CHECK: OpReturn
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Vertex %main "main"
+               OpSource HLSL 600
+               OpName %type__Globals "type.$Globals"
+               OpMemberName %type__Globals 0 "x"
+               OpMemberName %type__Globals 1 "y"
+               OpMemberName %type__Globals 2 "z"
+               OpName %_Globals "$Globals"
+               OpName %main "main"
+               OpDecorate %_Globals DescriptorSet 0
+               OpDecorate %_Globals Binding 0
+               OpMemberDecorate %type__Globals 0 Offset 0
+               OpMemberDecorate %type__Globals 1 Offset 4
+               OpMemberDecorate %type__Globals 2 Offset 8
+               OpDecorate %type__Globals Block
+      %float = OpTypeFloat 32
+%type__Globals = OpTypeStruct %float %float %float
+%_ptr_Uniform_type__Globals = OpTypePointer Uniform %type__Globals
+       %void = OpTypeVoid
+          %7 = OpTypeFunction %void
+   %_Globals = OpVariable %_ptr_Uniform_type__Globals Uniform
+       %main = OpFunction %void None %7
+          %8 = OpLabel
+          %9 = OpLoad %type__Globals %_Globals
+         %10 = OpCompositeExtract %float %9 1
+         %11 = OpCompositeInsert %type__Globals %10 %9 1
+               OpReturn
+               OpFunctionEnd
+
+)";
+
+  SinglePassRunAndMatch<opt::EliminateDeadMembersPass>(text, true);
+}
+
+TEST_F(EliminateDeadMemberTest, RemoveMembersUpdateInserExtract3) {
+  // Test that the members "x" and "z" are removed, and one member from the
+  // substruct. Update the OpCompositeExtract instruction. Update the
+  // OpCompositeInsert instruction.
+  const std::string text = R"(
+; CHECK: OpName
+; CHECK-NEXT: OpMemberName %type__Globals 0 "y"
+; CHECK-NOT: OpMemberName
+; CHECK: OpMemberDecorate %type__Globals 0 Offset 16
+; CHECK-NOT: OpMemberDecorate %type__Globals 1 Offset
+; CHECK: OpMemberDecorate [[struct:%\w+]] 0 Offset 4
+; CHECK: [[struct:%\w+]] = OpTypeStruct %float
+; CHECK: %type__Globals = OpTypeStruct [[struct]]
+; CHECK: [[ld:%\w+]] = OpLoad %type__Globals %_Globals
+; CHECK: [[ex:%\w+]] = OpCompositeExtract %float [[ld]] 0 0
+; CHECK: OpCompositeInsert %type__Globals [[ex]] [[ld]] 0 0
+; CHECK: OpReturn
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Vertex %main "main"
+               OpSource HLSL 600
+               OpName %type__Globals "type.$Globals"
+               OpMemberName %type__Globals 0 "x"
+               OpMemberName %type__Globals 1 "y"
+               OpMemberName %type__Globals 2 "z"
+               OpName %_Globals "$Globals"
+               OpName %main "main"
+               OpDecorate %_Globals DescriptorSet 0
+               OpDecorate %_Globals Binding 0
+               OpMemberDecorate %type__Globals 0 Offset 0
+               OpMemberDecorate %type__Globals 1 Offset 16
+               OpMemberDecorate %type__Globals 2 Offset 24
+               OpMemberDecorate %_struct_6 0 Offset 0
+               OpMemberDecorate %_struct_6 1 Offset 4
+               OpDecorate %type__Globals Block
+      %float = OpTypeFloat 32
+  %_struct_6 = OpTypeStruct %float %float
+%type__Globals = OpTypeStruct %float %_struct_6 %float
+%_ptr_Uniform_type__Globals = OpTypePointer Uniform %type__Globals
+       %void = OpTypeVoid
+          %7 = OpTypeFunction %void
+   %_Globals = OpVariable %_ptr_Uniform_type__Globals Uniform
+       %main = OpFunction %void None %7
+          %8 = OpLabel
+          %9 = OpLoad %type__Globals %_Globals
+         %10 = OpCompositeExtract %float %9 1 1
+         %11 = OpCompositeInsert %type__Globals %10 %9 1 1
+               OpReturn
+               OpFunctionEnd
+
+)";
+
+  SinglePassRunAndMatch<opt::EliminateDeadMembersPass>(text, true);
+}
+
+TEST_F(EliminateDeadMemberTest, RemoveMembersUpdateInserExtract4) {
+  // Test that the members "x" and "z" are removed, and one member from the
+  // substruct. Update the OpCompositeExtract instruction. Update the
+  // OpCompositeInsert instruction.
+  const std::string text = R"(
+; CHECK: OpName
+; CHECK-NEXT: OpMemberName %type__Globals 0 "y"
+; CHECK-NOT: OpMemberName
+; CHECK: OpMemberDecorate %type__Globals 0 Offset 16
+; CHECK-NOT: OpMemberDecorate %type__Globals 1 Offset
+; CHECK: OpMemberDecorate [[struct:%\w+]] 0 Offset 4
+; CHECK: [[struct:%\w+]] = OpTypeStruct %float
+; CHECK: [[array:%\w+]] = OpTypeArray [[struct]]
+; CHECK: %type__Globals = OpTypeStruct [[array]]
+; CHECK: [[ld:%\w+]] = OpLoad %type__Globals %_Globals
+; CHECK: [[ex:%\w+]] = OpCompositeExtract %float [[ld]] 0 1 0
+; CHECK: OpCompositeInsert %type__Globals [[ex]] [[ld]] 0 1 0
+; CHECK: OpReturn
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Vertex %main "main"
+               OpSource HLSL 600
+               OpName %type__Globals "type.$Globals"
+               OpMemberName %type__Globals 0 "x"
+               OpMemberName %type__Globals 1 "y"
+               OpMemberName %type__Globals 2 "z"
+               OpName %_Globals "$Globals"
+               OpName %main "main"
+               OpDecorate %_Globals DescriptorSet 0
+               OpDecorate %_Globals Binding 0
+               OpMemberDecorate %type__Globals 0 Offset 0
+               OpMemberDecorate %type__Globals 1 Offset 16
+               OpMemberDecorate %type__Globals 2 Offset 80
+               OpMemberDecorate %_struct_6 0 Offset 0
+               OpMemberDecorate %_struct_6 1 Offset 4
+               OpDecorate %array ArrayStride 16
+               OpDecorate %type__Globals Block
+       %uint = OpTypeInt 32 0                         ; 32-bit int, sign-less
+     %uint_4 = OpConstant %uint 4
+      %float = OpTypeFloat 32
+  %_struct_6 = OpTypeStruct %float %float
+  %array = OpTypeArray %_struct_6 %uint_4
+%type__Globals = OpTypeStruct %float %array %float
+%_ptr_Uniform_type__Globals = OpTypePointer Uniform %type__Globals
+       %void = OpTypeVoid
+          %7 = OpTypeFunction %void
+   %_Globals = OpVariable %_ptr_Uniform_type__Globals Uniform
+       %main = OpFunction %void None %7
+          %8 = OpLabel
+          %9 = OpLoad %type__Globals %_Globals
+         %10 = OpCompositeExtract %float %9 1 1 1
+         %11 = OpCompositeInsert %type__Globals %10 %9 1 1 1
+               OpReturn
+               OpFunctionEnd
+
+)";
+
+  SinglePassRunAndMatch<opt::EliminateDeadMembersPass>(text, true);
+}
+
+TEST_F(EliminateDeadMemberTest, RemoveMembersUpdateArrayLength) {
+  // Test that the members "x" and "y" are removed.
+  // Member "z" is live because of the OpArrayLength instruction.
+  // Update the OpArrayLength instruction.
+  const std::string text = R"(
+; CHECK: OpName
+; CHECK-NEXT: OpMemberName %type__Globals 0 "z"
+; CHECK-NOT: OpMemberName
+; CHECK: OpMemberDecorate %type__Globals 0 Offset 16
+; CHECK-NOT: OpMemberDecorate %type__Globals 1 Offset
+; CHECK: %type__Globals = OpTypeStruct %_runtimearr_float
+; CHECK: OpArrayLength %uint %_Globals 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Vertex %main "main"
+               OpSource HLSL 600
+               OpName %type__Globals "type.$Globals"
+               OpMemberName %type__Globals 0 "x"
+               OpMemberName %type__Globals 1 "y"
+               OpMemberName %type__Globals 2 "z"
+               OpName %_Globals "$Globals"
+               OpName %main "main"
+               OpDecorate %_Globals DescriptorSet 0
+               OpDecorate %_Globals Binding 0
+               OpMemberDecorate %type__Globals 0 Offset 0
+               OpMemberDecorate %type__Globals 1 Offset 4
+               OpMemberDecorate %type__Globals 2 Offset 16
+               OpDecorate %type__Globals Block
+       %uint = OpTypeInt 32 0
+      %float = OpTypeFloat 32
+%_runtimearr_float = OpTypeRuntimeArray %float
+%type__Globals = OpTypeStruct %float %float %_runtimearr_float
+%_ptr_Uniform_type__Globals = OpTypePointer Uniform %type__Globals
+       %void = OpTypeVoid
+          %9 = OpTypeFunction %void
+   %_Globals = OpVariable %_ptr_Uniform_type__Globals Uniform
+       %main = OpFunction %void None %9
+         %10 = OpLabel
+         %11 = OpLoad %type__Globals %_Globals
+         %12 = OpArrayLength %uint %_Globals 2
+               OpReturn
+               OpFunctionEnd
+)";
+
+  SinglePassRunAndMatch<opt::EliminateDeadMembersPass>(text, true);
+}
+
+TEST_F(EliminateDeadMemberTest, KeepMembersOpStore) {
+  // Test that all members are kept because of an OpStore.
+  // No change expected.
+  const std::string text = R"(
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Vertex %main "main"
+               OpSource HLSL 600
+               OpName %type__Globals "type.$Globals"
+               OpMemberName %type__Globals 0 "x"
+               OpMemberName %type__Globals 1 "y"
+               OpMemberName %type__Globals 2 "z"
+               OpName %_Globals "$Globals"
+               OpName %_Globals "$Globals2"
+               OpName %main "main"
+               OpDecorate %_Globals DescriptorSet 0
+               OpDecorate %_Globals Binding 0
+               OpMemberDecorate %type__Globals 0 Offset 0
+               OpMemberDecorate %type__Globals 1 Offset 4
+               OpMemberDecorate %type__Globals 2 Offset 16
+               OpDecorate %type__Globals Block
+       %uint = OpTypeInt 32 0
+      %float = OpTypeFloat 32
+%type__Globals = OpTypeStruct %float %float %float
+%_ptr_Uniform_type__Globals = OpTypePointer Uniform %type__Globals
+       %void = OpTypeVoid
+          %9 = OpTypeFunction %void
+   %_Globals = OpVariable %_ptr_Uniform_type__Globals Uniform
+   %_Globals2 = OpVariable %_ptr_Uniform_type__Globals Uniform
+       %main = OpFunction %void None %9
+         %10 = OpLabel
+         %11 = OpLoad %type__Globals %_Globals
+               OpStore %_Globals2 %11
+               OpReturn
+               OpFunctionEnd
+)";
+
+  auto result = SinglePassRunAndDisassemble<opt::EliminateDeadMembersPass>(
+      text, /* skip_nop = */ true, /* do_validation = */ true);
+  EXPECT_EQ(opt::Pass::Status::SuccessWithoutChange, std::get<1>(result));
+}
+
+TEST_F(EliminateDeadMemberTest, KeepMembersOpCopyMemory) {
+  // Test that all members are kept because of an OpCopyMemory.
+  // No change expected.
+  const std::string text = R"(
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Vertex %main "main"
+               OpSource HLSL 600
+               OpName %type__Globals "type.$Globals"
+               OpMemberName %type__Globals 0 "x"
+               OpMemberName %type__Globals 1 "y"
+               OpMemberName %type__Globals 2 "z"
+               OpName %_Globals "$Globals"
+               OpName %_Globals "$Globals2"
+               OpName %main "main"
+               OpDecorate %_Globals DescriptorSet 0
+               OpDecorate %_Globals Binding 0
+               OpMemberDecorate %type__Globals 0 Offset 0
+               OpMemberDecorate %type__Globals 1 Offset 4
+               OpMemberDecorate %type__Globals 2 Offset 16
+               OpDecorate %type__Globals Block
+       %uint = OpTypeInt 32 0
+      %float = OpTypeFloat 32
+%type__Globals = OpTypeStruct %float %float %float
+%_ptr_Uniform_type__Globals = OpTypePointer Uniform %type__Globals
+       %void = OpTypeVoid
+          %9 = OpTypeFunction %void
+   %_Globals = OpVariable %_ptr_Uniform_type__Globals Uniform
+   %_Globals2 = OpVariable %_ptr_Uniform_type__Globals Uniform
+       %main = OpFunction %void None %9
+         %10 = OpLabel
+               OpCopyMemory %_Globals2 %_Globals
+               OpReturn
+               OpFunctionEnd
+)";
+
+  auto result = SinglePassRunAndDisassemble<opt::EliminateDeadMembersPass>(
+      text, /* skip_nop = */ true, /* do_validation = */ true);
+  EXPECT_EQ(opt::Pass::Status::SuccessWithoutChange, std::get<1>(result));
+}
+
+TEST_F(EliminateDeadMemberTest, KeepMembersOpCopyMemorySized) {
+  // Test that all members are kept because of an OpCopyMemorySized.
+  // No change expected.
+  const std::string text = R"(
+               OpCapability Shader
+               OpCapability Addresses
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Vertex %main "main"
+               OpSource HLSL 600
+               OpName %type__Globals "type.$Globals"
+               OpMemberName %type__Globals 0 "x"
+               OpMemberName %type__Globals 1 "y"
+               OpMemberName %type__Globals 2 "z"
+               OpName %_Globals "$Globals"
+               OpName %_Globals "$Globals2"
+               OpName %main "main"
+               OpDecorate %_Globals DescriptorSet 0
+               OpDecorate %_Globals Binding 0
+               OpMemberDecorate %type__Globals 0 Offset 0
+               OpMemberDecorate %type__Globals 1 Offset 4
+               OpMemberDecorate %type__Globals 2 Offset 16
+               OpDecorate %type__Globals Block
+       %uint = OpTypeInt 32 0
+    %uint_20 = OpConstant %uint 20
+      %float = OpTypeFloat 32
+%type__Globals = OpTypeStruct %float %float %float
+%_ptr_Uniform_type__Globals = OpTypePointer Uniform %type__Globals
+       %void = OpTypeVoid
+          %9 = OpTypeFunction %void
+   %_Globals = OpVariable %_ptr_Uniform_type__Globals Uniform
+   %_Globals2 = OpVariable %_ptr_Uniform_type__Globals Uniform
+       %main = OpFunction %void None %9
+         %10 = OpLabel
+               OpCopyMemorySized %_Globals2 %_Globals %uint_20
+               OpReturn
+               OpFunctionEnd
+)";
+
+  auto result = SinglePassRunAndDisassemble<opt::EliminateDeadMembersPass>(
+      text, /* skip_nop = */ true, /* do_validation = */ true);
+  EXPECT_EQ(opt::Pass::Status::SuccessWithoutChange, std::get<1>(result));
+}
+
+TEST_F(EliminateDeadMemberTest, KeepMembersOpReturnValue) {
+  // Test that all members are kept because of an OpCopyMemorySized.
+  // No change expected.
+  const std::string text = R"(
+               OpCapability Shader
+               OpCapability Linkage
+               OpMemoryModel Logical GLSL450
+               OpSource HLSL 600
+               OpName %type__Globals "type.$Globals"
+               OpMemberName %type__Globals 0 "x"
+               OpMemberName %type__Globals 1 "y"
+               OpMemberName %type__Globals 2 "z"
+               OpName %_Globals "$Globals"
+               OpName %_Globals "$Globals2"
+               OpName %main "main"
+               OpDecorate %_Globals DescriptorSet 0
+               OpDecorate %_Globals Binding 0
+               OpMemberDecorate %type__Globals 0 Offset 0
+               OpMemberDecorate %type__Globals 1 Offset 4
+               OpMemberDecorate %type__Globals 2 Offset 16
+               OpDecorate %type__Globals Block
+       %uint = OpTypeInt 32 0
+    %uint_20 = OpConstant %uint 20
+      %float = OpTypeFloat 32
+%type__Globals = OpTypeStruct %float %float %float
+%_ptr_Uniform_type__Globals = OpTypePointer Uniform %type__Globals
+       %void = OpTypeVoid
+          %9 = OpTypeFunction %type__Globals
+   %_Globals = OpVariable %_ptr_Uniform_type__Globals Uniform
+   %_Globals2 = OpVariable %_ptr_Uniform_type__Globals Uniform
+       %main = OpFunction %type__Globals None %9
+         %10 = OpLabel
+         %11 = OpLoad %type__Globals %_Globals
+               OpReturnValue %11
+               OpFunctionEnd
+)";
+
+  auto result = SinglePassRunAndDisassemble<opt::EliminateDeadMembersPass>(
+      text, /* skip_nop = */ true, /* do_validation = */ true);
+  EXPECT_EQ(opt::Pass::Status::SuccessWithoutChange, std::get<1>(result));
+}
+
+TEST_F(EliminateDeadMemberTest, RemoveMemberAccessChainWithArrays) {
+  // Leave only 1 member in each of the structs.
+  // Update OpMemberName, OpMemberDecorate, and OpAccessChain.
+  const std::string text = R"(
+; CHECK: OpName
+; CHECK-NEXT: OpMemberName %type__Globals 0 "y"
+; CHECK-NOT: OpMemberName
+; CHECK: OpMemberDecorate %type__Globals 0 Offset 16
+; CHECK: OpMemberDecorate [[struct:%\w+]] 0 Offset 4
+; CHECK: [[struct]] = OpTypeStruct %float
+; CHECK: [[array:%\w+]] = OpTypeArray [[struct]]
+; CHECK: %type__Globals = OpTypeStruct [[array]]
+; CHECK: [[undef:%\w+]] = OpUndef %uint
+; CHECK: OpAccessChain %_ptr_Uniform_float %_Globals [[undef]] %uint_0 [[undef]] %uint_0
+               OpCapability Shader
+               OpCapability VariablePointersStorageBuffer
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Vertex %main "main"
+               OpSource HLSL 600
+               OpName %type__Globals "type.$Globals"
+               OpMemberName %type__Globals 0 "x"
+               OpMemberName %type__Globals 1 "y"
+               OpMemberName %type__Globals 2 "z"
+               OpName %_Globals "$Globals"
+               OpName %main "main"
+               OpDecorate %_Globals DescriptorSet 0
+               OpDecorate %_Globals Binding 0
+               OpMemberDecorate %type__Globals 0 Offset 0
+               OpMemberDecorate %type__Globals 1 Offset 16
+               OpMemberDecorate %type__Globals 2 Offset 48
+               OpMemberDecorate %_struct_4 0 Offset 0
+               OpMemberDecorate %_struct_4 1 Offset 4
+               OpDecorate %_arr__struct_4_uint_2 ArrayStride 16
+               OpDecorate %type__Globals Block
+       %uint = OpTypeInt 32 0
+     %uint_0 = OpConstant %uint 0
+     %uint_1 = OpConstant %uint 1
+     %uint_2 = OpConstant %uint 2
+     %uint_3 = OpConstant %uint 3
+      %float = OpTypeFloat 32
+  %_struct_4 = OpTypeStruct %float %float
+%_arr__struct_4_uint_2 = OpTypeArray %_struct_4 %uint_2
+%type__Globals = OpTypeStruct %float %_arr__struct_4_uint_2 %float
+%_arr_type__Globals_uint_3 = OpTypeArray %type__Globals %uint_3
+%_ptr_Uniform__arr_type__Globals_uint_3 = OpTypePointer Uniform %_arr_type__Globals_uint_3
+       %void = OpTypeVoid
+         %15 = OpTypeFunction %void
+%_ptr_Uniform_float = OpTypePointer Uniform %float
+   %_Globals = OpVariable %_ptr_Uniform__arr_type__Globals_uint_3 Uniform
+       %main = OpFunction %void None %15
+         %17 = OpLabel
+         %18 = OpUndef %uint
+         %19 = OpAccessChain %_ptr_Uniform_float %_Globals %18 %uint_1 %18 %uint_1
+               OpReturn
+               OpFunctionEnd
+)";
+
+  SinglePassRunAndMatch<opt::EliminateDeadMembersPass>(text, true);
+}
+
+TEST_F(EliminateDeadMemberTest, RemoveMemberInboundsAccessChain) {
+  // Test that the member "y" is removed.
+  // Update OpMemberName for |y| and |z|.
+  // Update OpMemberDecorate for |y| and |z|.
+  // Update OpInboundsAccessChain for access to |z|.
+  const std::string text = R"(
+; CHECK: OpName
+; CHECK-NEXT: OpMemberName %type__Globals 0 "x"
+; CHECK-NEXT: OpMemberName %type__Globals 1 "z"
+; CHECK-NOT: OpMemberName
+; CHECK: OpMemberDecorate %type__Globals 0 Offset 0
+; CHECK: OpMemberDecorate %type__Globals 1 Offset 8
+; CHECK: %type__Globals = OpTypeStruct %float %float
+; CHECK: OpInBoundsAccessChain %_ptr_Uniform_float %_Globals %int_0
+; CHECK: OpInBoundsAccessChain %_ptr_Uniform_float %_Globals %uint_1
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Vertex %main "main" %in_var_Position %gl_Position
+               OpSource HLSL 600
+               OpName %type__Globals "type.$Globals"
+               OpMemberName %type__Globals 0 "x"
+               OpMemberName %type__Globals 1 "y"
+               OpMemberName %type__Globals 2 "z"
+               OpName %_Globals "$Globals"
+               OpName %in_var_Position "in.var.Position"
+               OpName %main "main"
+               OpDecorate %gl_Position BuiltIn Position
+               OpDecorate %in_var_Position Location 0
+               OpDecorate %_Globals DescriptorSet 0
+               OpDecorate %_Globals Binding 0
+               OpMemberDecorate %type__Globals 0 Offset 0
+               OpMemberDecorate %type__Globals 1 Offset 4
+               OpMemberDecorate %type__Globals 2 Offset 8
+               OpDecorate %type__Globals Block
+        %int = OpTypeInt 32 1
+      %int_0 = OpConstant %int 0
+      %float = OpTypeFloat 32
+      %int_2 = OpConstant %int 2
+%type__Globals = OpTypeStruct %float %float %float
+%_ptr_Uniform_type__Globals = OpTypePointer Uniform %type__Globals
+    %v4float = OpTypeVector %float 4
+%_ptr_Input_v4float = OpTypePointer Input %v4float
+%_ptr_Output_v4float = OpTypePointer Output %v4float
+       %void = OpTypeVoid
+         %15 = OpTypeFunction %void
+%_ptr_Uniform_float = OpTypePointer Uniform %float
+   %_Globals = OpVariable %_ptr_Uniform_type__Globals Uniform
+%in_var_Position = OpVariable %_ptr_Input_v4float Input
+%gl_Position = OpVariable %_ptr_Output_v4float Output
+       %main = OpFunction %void None %15
+         %17 = OpLabel
+         %18 = OpLoad %v4float %in_var_Position
+         %19 = OpInBoundsAccessChain %_ptr_Uniform_float %_Globals %int_0
+         %20 = OpLoad %float %19
+         %21 = OpCompositeExtract %float %18 0
+         %22 = OpFAdd %float %21 %20
+         %23 = OpCompositeInsert %v4float %22 %18 0
+         %24 = OpCompositeExtract %float %18 1
+         %25 = OpCompositeInsert %v4float %24 %23 1
+         %26 = OpInBoundsAccessChain %_ptr_Uniform_float %_Globals %int_2
+         %27 = OpLoad %float %26
+         %28 = OpCompositeExtract %float %18 2
+         %29 = OpFAdd %float %28 %27
+         %30 = OpCompositeInsert %v4float %29 %25 2
+               OpStore %gl_Position %30
+               OpReturn
+               OpFunctionEnd
+)";
+
+  SinglePassRunAndMatch<opt::EliminateDeadMembersPass>(text, true);
+}
+
+TEST_F(EliminateDeadMemberTest, RemoveMemberPtrAccessChain) {
+  // Test that the member "y" is removed.
+  // Update OpMemberName for |y| and |z|.
+  // Update OpMemberDecorate for |y| and |z|.
+  // Update OpInboundsAccessChain for access to |z|.
+  const std::string text = R"(
+; CHECK: OpName
+; CHECK-NEXT: OpMemberName %type__Globals 0 "x"
+; CHECK-NEXT: OpMemberName %type__Globals 1 "z"
+; CHECK-NOT: OpMemberName
+; CHECK: OpMemberDecorate %type__Globals 0 Offset 0
+; CHECK: OpMemberDecorate %type__Globals 1 Offset 16
+; CHECK: %type__Globals = OpTypeStruct %float %float
+; CHECK: [[ac:%\w+]] = OpAccessChain %_ptr_Uniform_type__Globals %_Globals %uint_0
+; CHECK: OpPtrAccessChain %_ptr_Uniform_float [[ac]] %uint_1 %uint_0
+; CHECK: OpPtrAccessChain %_ptr_Uniform_float [[ac]] %uint_0 %uint_1
+               OpCapability Shader
+               OpCapability VariablePointersStorageBuffer
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Vertex %main "main"
+               OpSource HLSL 600
+               OpName %type__Globals "type.$Globals"
+               OpMemberName %type__Globals 0 "x"
+               OpMemberName %type__Globals 1 "y"
+               OpMemberName %type__Globals 2 "z"
+               OpName %_Globals "$Globals"
+               OpName %main "main"
+               OpDecorate %_Globals DescriptorSet 0
+               OpDecorate %_Globals Binding 0
+               OpMemberDecorate %type__Globals 0 Offset 0
+               OpMemberDecorate %type__Globals 1 Offset 4
+               OpMemberDecorate %type__Globals 2 Offset 16
+               OpDecorate %type__Globals Block
+       %uint = OpTypeInt 32 0
+     %uint_0 = OpConstant %uint 0
+     %uint_1 = OpConstant %uint 1
+     %uint_2 = OpConstant %uint 2
+     %uint_3 = OpConstant %uint 3
+      %float = OpTypeFloat 32
+%type__Globals = OpTypeStruct %float %float %float
+%_arr_type__Globals_uint_3 = OpTypeArray %type__Globals %uint_3
+%_ptr_Uniform_type__Globals = OpTypePointer Uniform %type__Globals
+%_ptr_Uniform__arr_type__Globals_uint_3 = OpTypePointer Uniform %_arr_type__Globals_uint_3
+       %void = OpTypeVoid
+         %14 = OpTypeFunction %void
+%_ptr_Uniform_float = OpTypePointer Uniform %float
+   %_Globals = OpVariable %_ptr_Uniform__arr_type__Globals_uint_3 Uniform
+       %main = OpFunction %void None %14
+         %16 = OpLabel
+         %17 = OpAccessChain %_ptr_Uniform_type__Globals %_Globals %uint_0
+         %18 = OpPtrAccessChain %_ptr_Uniform_float %17 %uint_1 %uint_0
+         %19 = OpPtrAccessChain %_ptr_Uniform_float %17 %uint_0 %uint_2
+               OpReturn
+               OpFunctionEnd
+)";
+
+  SinglePassRunAndMatch<opt::EliminateDeadMembersPass>(text, true);
+}
+
+TEST_F(EliminateDeadMemberTest, RemoveMemberInBoundsPtrAccessChain) {
+  // Test that the member "y" is removed.
+  // Update OpMemberName for |y| and |z|.
+  // Update OpMemberDecorate for |y| and |z|.
+  // Update OpInboundsAccessChain for access to |z|.
+  const std::string text = R"(
+; CHECK: OpName
+; CHECK-NEXT: OpMemberName %type__Globals 0 "x"
+; CHECK-NEXT: OpMemberName %type__Globals 1 "z"
+; CHECK-NOT: OpMemberName
+; CHECK: OpMemberDecorate %type__Globals 0 Offset 0
+; CHECK: OpMemberDecorate %type__Globals 1 Offset 16
+; CHECK: %type__Globals = OpTypeStruct %float %float
+; CHECK: [[ac:%\w+]] = OpAccessChain %_ptr_Uniform_type__Globals %_Globals %uint_0
+; CHECK: OpInBoundsPtrAccessChain %_ptr_Uniform_float [[ac]] %uint_1 %uint_0
+; CHECK: OpInBoundsPtrAccessChain %_ptr_Uniform_float [[ac]] %uint_0 %uint_1
+               OpCapability Shader
+               OpCapability Addresses
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Vertex %main "main"
+               OpSource HLSL 600
+               OpName %type__Globals "type.$Globals"
+               OpMemberName %type__Globals 0 "x"
+               OpMemberName %type__Globals 1 "y"
+               OpMemberName %type__Globals 2 "z"
+               OpName %_Globals "$Globals"
+               OpName %main "main"
+               OpDecorate %_Globals DescriptorSet 0
+               OpDecorate %_Globals Binding 0
+               OpMemberDecorate %type__Globals 0 Offset 0
+               OpMemberDecorate %type__Globals 1 Offset 4
+               OpMemberDecorate %type__Globals 2 Offset 16
+               OpDecorate %type__Globals Block
+       %uint = OpTypeInt 32 0
+     %uint_0 = OpConstant %uint 0
+     %uint_1 = OpConstant %uint 1
+     %uint_2 = OpConstant %uint 2
+     %uint_3 = OpConstant %uint 3
+      %float = OpTypeFloat 32
+%type__Globals = OpTypeStruct %float %float %float
+%_arr_type__Globals_uint_3 = OpTypeArray %type__Globals %uint_3
+%_ptr_Uniform_type__Globals = OpTypePointer Uniform %type__Globals
+%_ptr_Uniform__arr_type__Globals_uint_3 = OpTypePointer Uniform %_arr_type__Globals_uint_3
+       %void = OpTypeVoid
+         %14 = OpTypeFunction %void
+%_ptr_Uniform_float = OpTypePointer Uniform %float
+   %_Globals = OpVariable %_ptr_Uniform__arr_type__Globals_uint_3 Uniform
+       %main = OpFunction %void None %14
+         %16 = OpLabel
+         %17 = OpAccessChain %_ptr_Uniform_type__Globals %_Globals %uint_0
+         %18 = OpInBoundsPtrAccessChain %_ptr_Uniform_float %17 %uint_1 %uint_0
+         %19 = OpInBoundsPtrAccessChain %_ptr_Uniform_float %17 %uint_0 %uint_2
+               OpReturn
+               OpFunctionEnd
+)";
+
+  SinglePassRunAndMatch<opt::EliminateDeadMembersPass>(text, true);
+}
+
+TEST_F(EliminateDeadMemberTest, DontRemoveModfStructResultTypeMembers) {
+  const std::string text = R"(
+               OpCapability Shader
+          %1 = OpExtInstImport "GLSL.std.450"
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Fragment %PS_TerrainElevation "PS_TerrainElevation"
+               OpExecutionMode %PS_TerrainElevation OriginUpperLeft
+               OpSource HLSL 600
+      %float = OpTypeFloat 32
+       %void = OpTypeVoid
+         %21 = OpTypeFunction %void
+%ModfStructType = OpTypeStruct %float %float
+%PS_TerrainElevation = OpFunction %void None %21
+         %22 = OpLabel
+         %23 = OpUndef %float
+         %24 = OpExtInst %ModfStructType %1 ModfStruct %23
+         %25 = OpCompositeExtract %float %24 1
+               OpReturn
+               OpFunctionEnd
+)";
+
+  auto result = SinglePassRunAndDisassemble<opt::EliminateDeadMembersPass>(
+      text, /* skip_nop = */ true, /* do_validation = */ true);
+  EXPECT_EQ(opt::Pass::Status::SuccessWithoutChange, std::get<1>(result));
+}
+
+}  // namespace

3rdparty/spirv-tools/test/opt/inst_bindless_check_test.cpp

@@ -1580,6 +1580,329 @@ OpFunctionEnd
       true);
 }
 
+TEST_F(InstBindlessTest, MultipleDebugFunctions) {
+  // Same source as Simple, but compiled -g and not optimized, especially not
+  // inlined. The OpSource has had the source extracted for the sake of brevity.
+
+  const std::string defs_before =
+      R"(OpCapability Shader
+%2 = OpExtInstImport "GLSL.std.450"
+OpMemoryModel Logical GLSL450
+OpEntryPoint Fragment %MainPs "MainPs" %i_vTextureCoords %_entryPointOutput_vColor
+OpExecutionMode %MainPs OriginUpperLeft
+%1 = OpString "foo5.frag"
+OpSource HLSL 500 %1
+OpName %MainPs "MainPs"
+OpName %PS_INPUT "PS_INPUT"
+OpMemberName %PS_INPUT 0 "vTextureCoords"
+OpName %PS_OUTPUT "PS_OUTPUT"
+OpMemberName %PS_OUTPUT 0 "vColor"
+OpName %_MainPs_struct_PS_INPUT_vf21_ "@MainPs(struct-PS_INPUT-vf21;"
+OpName %i "i"
+OpName %ps_output "ps_output"
+OpName %g_tColor "g_tColor"
+OpName %PerViewConstantBuffer_t "PerViewConstantBuffer_t"
+OpMemberName %PerViewConstantBuffer_t 0 "g_nDataIdx"
+OpName %_ ""
+OpName %g_sAniso "g_sAniso"
+OpName %i_0 "i"
+OpName %i_vTextureCoords "i.vTextureCoords"
+OpName %_entryPointOutput_vColor "@entryPointOutput.vColor"
+OpName %param "param"
+OpDecorate %g_tColor DescriptorSet 0
+OpDecorate %g_tColor Binding 0
+OpMemberDecorate %PerViewConstantBuffer_t 0 Offset 0
+OpDecorate %PerViewConstantBuffer_t Block
+OpDecorate %g_sAniso DescriptorSet 0
+OpDecorate %g_sAniso Binding 1
+OpDecorate %i_vTextureCoords Location 0
+OpDecorate %_entryPointOutput_vColor Location 0
+%void = OpTypeVoid
+%4 = OpTypeFunction %void
+%float = OpTypeFloat 32
+%v2float = OpTypeVector %float 2
+%PS_INPUT = OpTypeStruct %v2float
+%_ptr_Function_PS_INPUT = OpTypePointer Function %PS_INPUT
+%v4float = OpTypeVector %float 4
+%PS_OUTPUT = OpTypeStruct %v4float
+%13 = OpTypeFunction %PS_OUTPUT %_ptr_Function_PS_INPUT
+%_ptr_Function_PS_OUTPUT = OpTypePointer Function %PS_OUTPUT
+%int = OpTypeInt 32 1
+%int_0 = OpConstant %int 0
+%21 = OpTypeImage %float 2D 0 0 0 1 Unknown
+%uint = OpTypeInt 32 0
+%uint_128 = OpConstant %uint 128
+%_arr_21_uint_128 = OpTypeArray %21 %uint_128
+%_ptr_UniformConstant__arr_21_uint_128 = OpTypePointer UniformConstant %_arr_21_uint_128
+%g_tColor = OpVariable %_ptr_UniformConstant__arr_21_uint_128 UniformConstant
+%PerViewConstantBuffer_t = OpTypeStruct %uint
+%_ptr_PushConstant_PerViewConstantBuffer_t = OpTypePointer PushConstant %PerViewConstantBuffer_t
+%_ = OpVariable %_ptr_PushConstant_PerViewConstantBuffer_t PushConstant
+%_ptr_PushConstant_uint = OpTypePointer PushConstant %uint
+%_ptr_UniformConstant_21 = OpTypePointer UniformConstant %21
+%36 = OpTypeSampler
+%_ptr_UniformConstant_36 = OpTypePointer UniformConstant %36
+%g_sAniso = OpVariable %_ptr_UniformConstant_36 UniformConstant
+%40 = OpTypeSampledImage %21
+%_ptr_Function_v2float = OpTypePointer Function %v2float
+%_ptr_Function_v4float = OpTypePointer Function %v4float
+%_ptr_Input_v2float = OpTypePointer Input %v2float
+%i_vTextureCoords = OpVariable %_ptr_Input_v2float Input
+%_ptr_Output_v4float = OpTypePointer Output %v4float
+%_entryPointOutput_vColor = OpVariable %_ptr_Output_v4float Output
+)";
+
+  const std::string defs_after =
+      R"(OpCapability Shader
+OpExtension "SPV_KHR_storage_buffer_storage_class"
+%1 = OpExtInstImport "GLSL.std.450"
+OpMemoryModel Logical GLSL450
+OpEntryPoint Fragment %MainPs "MainPs" %i_vTextureCoords %_entryPointOutput_vColor %gl_FragCoord
+OpExecutionMode %MainPs OriginUpperLeft
+%5 = OpString "foo5.frag"
+OpSource HLSL 500 %5
+OpName %MainPs "MainPs"
+OpName %PS_INPUT "PS_INPUT"
+OpMemberName %PS_INPUT 0 "vTextureCoords"
+OpName %PS_OUTPUT "PS_OUTPUT"
+OpMemberName %PS_OUTPUT 0 "vColor"
+OpName %_MainPs_struct_PS_INPUT_vf21_ "@MainPs(struct-PS_INPUT-vf21;"
+OpName %i "i"
+OpName %ps_output "ps_output"
+OpName %g_tColor "g_tColor"
+OpName %PerViewConstantBuffer_t "PerViewConstantBuffer_t"
+OpMemberName %PerViewConstantBuffer_t 0 "g_nDataIdx"
+OpName %_ ""
+OpName %g_sAniso "g_sAniso"
+OpName %i_0 "i"
+OpName %i_vTextureCoords "i.vTextureCoords"
+OpName %_entryPointOutput_vColor "@entryPointOutput.vColor"
+OpName %param "param"
+OpDecorate %g_tColor DescriptorSet 0
+OpDecorate %g_tColor Binding 0
+OpMemberDecorate %PerViewConstantBuffer_t 0 Offset 0
+OpDecorate %PerViewConstantBuffer_t Block
+OpDecorate %g_sAniso DescriptorSet 0
+OpDecorate %g_sAniso Binding 1
+OpDecorate %i_vTextureCoords Location 0
+OpDecorate %_entryPointOutput_vColor Location 0
+OpDecorate %_runtimearr_uint ArrayStride 4
+OpDecorate %_struct_77 Block
+OpMemberDecorate %_struct_77 0 Offset 0
+OpMemberDecorate %_struct_77 1 Offset 4
+OpDecorate %79 DescriptorSet 7
+OpDecorate %79 Binding 0
+OpDecorate %gl_FragCoord BuiltIn FragCoord
+%void = OpTypeVoid
+%18 = OpTypeFunction %void
+%float = OpTypeFloat 32
+%v2float = OpTypeVector %float 2
+%PS_INPUT = OpTypeStruct %v2float
+%_ptr_Function_PS_INPUT = OpTypePointer Function %PS_INPUT
+%v4float = OpTypeVector %float 4
+%PS_OUTPUT = OpTypeStruct %v4float
+%23 = OpTypeFunction %PS_OUTPUT %_ptr_Function_PS_INPUT
+%_ptr_Function_PS_OUTPUT = OpTypePointer Function %PS_OUTPUT
+%int = OpTypeInt 32 1
+%int_0 = OpConstant %int 0
+%27 = OpTypeImage %float 2D 0 0 0 1 Unknown
+%uint = OpTypeInt 32 0
+%uint_128 = OpConstant %uint 128
+%_arr_27_uint_128 = OpTypeArray %27 %uint_128
+%_ptr_UniformConstant__arr_27_uint_128 = OpTypePointer UniformConstant %_arr_27_uint_128
+%g_tColor = OpVariable %_ptr_UniformConstant__arr_27_uint_128 UniformConstant
+%PerViewConstantBuffer_t = OpTypeStruct %uint
+%_ptr_PushConstant_PerViewConstantBuffer_t = OpTypePointer PushConstant %PerViewConstantBuffer_t
+%_ = OpVariable %_ptr_PushConstant_PerViewConstantBuffer_t PushConstant
+%_ptr_PushConstant_uint = OpTypePointer PushConstant %uint
+%_ptr_UniformConstant_27 = OpTypePointer UniformConstant %27
+%35 = OpTypeSampler
+%_ptr_UniformConstant_35 = OpTypePointer UniformConstant %35
+%g_sAniso = OpVariable %_ptr_UniformConstant_35 UniformConstant
+%37 = OpTypeSampledImage %27
+%_ptr_Function_v2float = OpTypePointer Function %v2float
+%_ptr_Function_v4float = OpTypePointer Function %v4float
+%_ptr_Input_v2float = OpTypePointer Input %v2float
+%i_vTextureCoords = OpVariable %_ptr_Input_v2float Input
+%_ptr_Output_v4float = OpTypePointer Output %v4float
+%_entryPointOutput_vColor = OpVariable %_ptr_Output_v4float Output
+%uint_0 = OpConstant %uint 0
+%bool = OpTypeBool
+%70 = OpTypeFunction %void %uint %uint %uint %uint
+%_runtimearr_uint = OpTypeRuntimeArray %uint
+%_struct_77 = OpTypeStruct %uint %_runtimearr_uint
+%_ptr_StorageBuffer__struct_77 = OpTypePointer StorageBuffer %_struct_77
+%79 = OpVariable %_ptr_StorageBuffer__struct_77 StorageBuffer
+%_ptr_StorageBuffer_uint = OpTypePointer StorageBuffer %uint
+%uint_9 = OpConstant %uint 9
+%uint_4 = OpConstant %uint 4
+%uint_1 = OpConstant %uint 1
+%uint_23 = OpConstant %uint 23
+%uint_2 = OpConstant %uint 2
+%uint_3 = OpConstant %uint 3
+%_ptr_Input_v4float = OpTypePointer Input %v4float
+%gl_FragCoord = OpVariable %_ptr_Input_v4float Input
+%v4uint = OpTypeVector %uint 4
+%uint_5 = OpConstant %uint 5
+%uint_6 = OpConstant %uint 6
+%uint_7 = OpConstant %uint 7
+%uint_8 = OpConstant %uint 8
+%uint_93 = OpConstant %uint 93
+%125 = OpConstantNull %v4float
+)";
+
+  const std::string func1_before =
+      R"(%MainPs = OpFunction %void None %4
+%6 = OpLabel
+%i_0 = OpVariable %_ptr_Function_PS_INPUT Function
+%param = OpVariable %_ptr_Function_PS_INPUT Function
+OpLine %1 21 0
+%54 = OpLoad %v2float %i_vTextureCoords
+%55 = OpAccessChain %_ptr_Function_v2float %i_0 %int_0
+OpStore %55 %54
+%59 = OpLoad %PS_INPUT %i_0
+OpStore %param %59
+%60 = OpFunctionCall %PS_OUTPUT %_MainPs_struct_PS_INPUT_vf21_ %param
+%61 = OpCompositeExtract %v4float %60 0
+OpStore %_entryPointOutput_vColor %61
+OpReturn
+OpFunctionEnd
+)";
+
+  const std::string func1_after =
+      R"(%MainPs = OpFunction %void None %18
+%42 = OpLabel
+%i_0 = OpVariable %_ptr_Function_PS_INPUT Function
+%param = OpVariable %_ptr_Function_PS_INPUT Function
+OpLine %5 21 0
+%43 = OpLoad %v2float %i_vTextureCoords
+%44 = OpAccessChain %_ptr_Function_v2float %i_0 %int_0
+OpStore %44 %43
+%45 = OpLoad %PS_INPUT %i_0
+OpStore %param %45
+%46 = OpFunctionCall %PS_OUTPUT %_MainPs_struct_PS_INPUT_vf21_ %param
+%47 = OpCompositeExtract %v4float %46 0
+OpStore %_entryPointOutput_vColor %47
+OpReturn
+OpFunctionEnd
+)";
+
+  const std::string func2_before =
+      R"(%_MainPs_struct_PS_INPUT_vf21_ = OpFunction %PS_OUTPUT None %13
+%i = OpFunctionParameter %_ptr_Function_PS_INPUT
+%16 = OpLabel
+%ps_output = OpVariable %_ptr_Function_PS_OUTPUT Function
+OpLine %1 24 0
+%31 = OpAccessChain %_ptr_PushConstant_uint %_ %int_0
+%32 = OpLoad %uint %31
+%34 = OpAccessChain %_ptr_UniformConstant_21 %g_tColor %32
+%35 = OpLoad %21 %34
+%39 = OpLoad %36 %g_sAniso
+%41 = OpSampledImage %40 %35 %39
+%43 = OpAccessChain %_ptr_Function_v2float %i %int_0
+%44 = OpLoad %v2float %43
+%45 = OpImageSampleImplicitLod %v4float %41 %44
+%47 = OpAccessChain %_ptr_Function_v4float %ps_output %int_0
+OpStore %47 %45
+OpLine %1 25 0
+%48 = OpLoad %PS_OUTPUT %ps_output
+OpReturnValue %48
+OpFunctionEnd
+)";
+
+  const std::string func2_after =
+      R"(%_MainPs_struct_PS_INPUT_vf21_ = OpFunction %PS_OUTPUT None %23
+%i = OpFunctionParameter %_ptr_Function_PS_INPUT
+%48 = OpLabel
+%ps_output = OpVariable %_ptr_Function_PS_OUTPUT Function
+OpLine %5 24 0
+%49 = OpAccessChain %_ptr_PushConstant_uint %_ %int_0
+%50 = OpLoad %uint %49
+%51 = OpAccessChain %_ptr_UniformConstant_27 %g_tColor %50
+%52 = OpLoad %27 %51
+%53 = OpLoad %35 %g_sAniso
+%54 = OpSampledImage %37 %52 %53
+%55 = OpAccessChain %_ptr_Function_v2float %i %int_0
+%56 = OpLoad %v2float %55
+%62 = OpULessThan %bool %50 %uint_128
+OpSelectionMerge %63 None
+OpBranchConditional %62 %64 %65
+%64 = OpLabel
+%66 = OpLoad %27 %51
+%67 = OpSampledImage %37 %66 %53
+%68 = OpImageSampleImplicitLod %v4float %67 %56
+OpBranch %63
+%65 = OpLabel
+%124 = OpFunctionCall %void %69 %uint_93 %uint_0 %50 %uint_128
+OpBranch %63
+%63 = OpLabel
+%126 = OpPhi %v4float %68 %64 %125 %65
+%58 = OpAccessChain %_ptr_Function_v4float %ps_output %int_0
+OpStore %58 %126
+OpLine %5 25 0
+%59 = OpLoad %PS_OUTPUT %ps_output
+OpReturnValue %59
+OpFunctionEnd
+)";
+
+  const std::string output_func =
+      R"(%69 = OpFunction %void None %70
+%71 = OpFunctionParameter %uint
+%72 = OpFunctionParameter %uint
+%73 = OpFunctionParameter %uint
+%74 = OpFunctionParameter %uint
+%75 = OpLabel
+%81 = OpAccessChain %_ptr_StorageBuffer_uint %79 %uint_0
+%84 = OpAtomicIAdd %uint %81 %uint_4 %uint_0 %uint_9
+%85 = OpIAdd %uint %84 %uint_9
+%86 = OpArrayLength %uint %79 1
+%87 = OpULessThanEqual %bool %85 %86
+OpSelectionMerge %88 None
+OpBranchConditional %87 %89 %88
+%89 = OpLabel
+%90 = OpIAdd %uint %84 %uint_0
+%92 = OpAccessChain %_ptr_StorageBuffer_uint %79 %uint_1 %90
+OpStore %92 %uint_9
+%94 = OpIAdd %uint %84 %uint_1
+%95 = OpAccessChain %_ptr_StorageBuffer_uint %79 %uint_1 %94
+OpStore %95 %uint_23
+%97 = OpIAdd %uint %84 %uint_2
+%98 = OpAccessChain %_ptr_StorageBuffer_uint %79 %uint_1 %97
+OpStore %98 %71
+%100 = OpIAdd %uint %84 %uint_3
+%101 = OpAccessChain %_ptr_StorageBuffer_uint %79 %uint_1 %100
+OpStore %101 %uint_4
+%104 = OpLoad %v4float %gl_FragCoord
+%106 = OpBitcast %v4uint %104
+%107 = OpCompositeExtract %uint %106 0
+%108 = OpIAdd %uint %84 %uint_4
+%109 = OpAccessChain %_ptr_StorageBuffer_uint %79 %uint_1 %108
+OpStore %109 %107
+%110 = OpCompositeExtract %uint %106 1
+%112 = OpIAdd %uint %84 %uint_5
+%113 = OpAccessChain %_ptr_StorageBuffer_uint %79 %uint_1 %112
+OpStore %113 %110
+%115 = OpIAdd %uint %84 %uint_6
+%116 = OpAccessChain %_ptr_StorageBuffer_uint %79 %uint_1 %115
+OpStore %116 %72
+%118 = OpIAdd %uint %84 %uint_7
+%119 = OpAccessChain %_ptr_StorageBuffer_uint %79 %uint_1 %118
+OpStore %119 %73
+%121 = OpIAdd %uint %84 %uint_8
+%122 = OpAccessChain %_ptr_StorageBuffer_uint %79 %uint_1 %121
+OpStore %122 %74
+OpBranch %88
+%88 = OpLabel
+OpReturn
+OpFunctionEnd
+)";
+
+  // SetAssembleOptions(SPV_TEXT_TO_BINARY_OPTION_PRESERVE_NUMERIC_IDS);
+  SinglePassRunAndCheck<InstBindlessCheckPass>(
+      defs_before + func1_before + func2_before,
+      defs_after + func1_after + func2_after + output_func, true, true);
+}
+
 TEST_F(InstBindlessTest, RuntimeArray) {
   // This test verifies that the pass will correctly instrument shader
   // with runtime descriptor array. This test was created by editing the

3rdparty/spirv-tools/test/opt/optimizer_test.cpp

@@ -242,112 +242,111 @@ TEST(Optimizer, WebGPUModeSetsCorrectPasses) {
     EXPECT_EQ(registered_passes[i], expected_passes[i]);
 }
 
-TEST(Optimizer, WebGPUModeFlattenDecorationsRuns) {
-  const std::string input = R"(OpCapability Shader
-OpCapability VulkanMemoryModelKHR
-OpExtension "SPV_KHR_vulkan_memory_model"
-OpMemoryModel Logical VulkanKHR
-OpEntryPoint Fragment %main "main" %hue %saturation %value
-OpExecutionMode %main OriginUpperLeft
-OpDecorate %group Flat
-OpDecorate %group NoPerspective
-%group = OpDecorationGroup
-%void = OpTypeVoid
-%void_fn = OpTypeFunction %void
-%float = OpTypeFloat 32
-%_ptr_Input_float = OpTypePointer Input %float
-%hue = OpVariable %_ptr_Input_float Input
-%saturation = OpVariable %_ptr_Input_float Input
-%value = OpVariable %_ptr_Input_float Input
-%main = OpFunction %void None %void_fn
-%entry = OpLabel
-OpReturn
-OpFunctionEnd
-)";
-
-  const std::string expected = R"(OpCapability Shader
-OpCapability VulkanMemoryModelKHR
-OpExtension "SPV_KHR_vulkan_memory_model"
-OpMemoryModel Logical VulkanKHR
-OpEntryPoint Fragment %1 "main" %2 %3 %4
-OpExecutionMode %1 OriginUpperLeft
-%void = OpTypeVoid
-%7 = OpTypeFunction %void
-%float = OpTypeFloat 32
-%_ptr_Input_float = OpTypePointer Input %float
-%2 = OpVariable %_ptr_Input_float Input
-%3 = OpVariable %_ptr_Input_float Input
-%4 = OpVariable %_ptr_Input_float Input
-%1 = OpFunction %void None %7
-%10 = OpLabel
-OpReturn
-OpFunctionEnd
-)";
-
+struct WebGPUPassCase {
+  // Input SPIR-V
+  std::string input;
+  // Expected result SPIR-V
+  std::string expected;
+  // Specific pass under test, used for logging messages.
+  std::string pass;
+};
+
+using WebGPUPassTest = PassTest<::testing::TestWithParam<WebGPUPassCase>>;
+
+TEST_P(WebGPUPassTest, Ran) {
   SpirvTools tools(SPV_ENV_WEBGPU_0);
   std::vector<uint32_t> binary;
-  tools.Assemble(input, &binary);
+  tools.Assemble(GetParam().input, &binary);
 
   Optimizer opt(SPV_ENV_WEBGPU_0);
   opt.RegisterWebGPUPasses();
 
   std::vector<uint32_t> optimized;
-  ValidatorOptions validator_options;
+  class ValidatorOptions validator_options;
   ASSERT_TRUE(opt.Run(binary.data(), binary.size(), &optimized,
                       validator_options, true));
 
   std::string disassembly;
   tools.Disassemble(optimized.data(), optimized.size(), &disassembly);
 
-  EXPECT_EQ(expected, disassembly);
+  EXPECT_EQ(GetParam().expected, disassembly)
+      << "Was expecting pass '" << GetParam().pass << "' to have been run.\n";
 }
 
-TEST(Optimizer, WebGPUModeStripDebugRuns) {
-  const std::string input = R"(OpCapability Shader
-OpCapability VulkanMemoryModelKHR
-OpExtension "SPV_KHR_vulkan_memory_model"
-OpMemoryModel Logical VulkanKHR
-OpEntryPoint Vertex %func "shader"
-OpName %main "main"
-OpName %void_fn "void_fn"
-%void = OpTypeVoid
-%void_f = OpTypeFunction %void
-%func = OpFunction %void None %void_f
-%label = OpLabel
-OpReturn
-OpFunctionEnd
-)";
-
-  const std::string expected = R"(OpCapability Shader
-OpCapability VulkanMemoryModelKHR
-OpExtension "SPV_KHR_vulkan_memory_model"
-OpMemoryModel Logical VulkanKHR
-OpEntryPoint Vertex %1 "shader"
-%void = OpTypeVoid
-%5 = OpTypeFunction %void
-%1 = OpFunction %void None %5
-%6 = OpLabel
-OpReturn
-OpFunctionEnd
-)";
-
-  SpirvTools tools(SPV_ENV_WEBGPU_0);
-  std::vector<uint32_t> binary;
-  tools.Assemble(input, &binary);
-
-  Optimizer opt(SPV_ENV_WEBGPU_0);
-  opt.RegisterWebGPUPasses();
-
-  std::vector<uint32_t> optimized;
-  ValidatorOptions validator_options;
-  ASSERT_TRUE(opt.Run(binary.data(), binary.size(), &optimized,
-                      validator_options, true));
-
-  std::string disassembly;
-  tools.Disassemble(optimized.data(), optimized.size(), &disassembly);
-
-  EXPECT_EQ(expected, disassembly);
-}
+INSTANTIATE_TEST_SUITE_P(
+    WebGPU, WebGPUPassTest,
+    ::testing::ValuesIn(std::vector<WebGPUPassCase>{
+        // FlattenDecorations
+        {// input
+         "OpCapability Shader\n"
+         "OpCapability VulkanMemoryModelKHR\n"
+         "OpExtension \"SPV_KHR_vulkan_memory_model\"\n"
+         "OpMemoryModel Logical VulkanKHR\n"
+         "OpEntryPoint Fragment %main \"main\" %hue %saturation %value\n"
+         "OpExecutionMode %main OriginUpperLeft\n"
+         "OpDecorate %group Flat\n"
+         "OpDecorate %group NoPerspective\n"
+         "%group = OpDecorationGroup\n"
+         "%void = OpTypeVoid\n"
+         "%void_fn = OpTypeFunction %void\n"
+         "%float = OpTypeFloat 32\n"
+         "%_ptr_Input_float = OpTypePointer Input %float\n"
+         "%hue = OpVariable %_ptr_Input_float Input\n"
+         "%saturation = OpVariable %_ptr_Input_float Input\n"
+         "%value = OpVariable %_ptr_Input_float Input\n"
+         "%main = OpFunction %void None %void_fn\n"
+         "%entry = OpLabel\n"
+         "OpReturn\n"
+         "OpFunctionEnd\n",
+         // expected
+         "OpCapability Shader\n"
+         "OpCapability VulkanMemoryModelKHR\n"
+         "OpExtension \"SPV_KHR_vulkan_memory_model\"\n"
+         "OpMemoryModel Logical VulkanKHR\n"
+         "OpEntryPoint Fragment %1 \"main\" %2 %3 %4\n"
+         "OpExecutionMode %1 OriginUpperLeft\n"
+         "%void = OpTypeVoid\n"
+         "%7 = OpTypeFunction %void\n"
+         "%float = OpTypeFloat 32\n"
+         "%_ptr_Input_float = OpTypePointer Input %float\n"
+         "%2 = OpVariable %_ptr_Input_float Input\n"
+         "%3 = OpVariable %_ptr_Input_float Input\n"
+         "%4 = OpVariable %_ptr_Input_float Input\n"
+         "%1 = OpFunction %void None %7\n"
+         "%10 = OpLabel\n"
+         "OpReturn\n"
+         "OpFunctionEnd\n",
+         // pass
+         "flatten-decorations"},
+        // Strip Debug
+        {// input
+         "OpCapability Shader\n"
+         "OpCapability VulkanMemoryModelKHR\n"
+         "OpExtension \"SPV_KHR_vulkan_memory_model\"\n"
+         "OpMemoryModel Logical VulkanKHR\n"
+         "OpEntryPoint Vertex %func \"shader\"\n"
+         "OpName %main \"main\"\n"
+         "OpName %void_fn \"void_fn\"\n"
+         "%void = OpTypeVoid\n"
+         "%void_f = OpTypeFunction %void\n"
+         "%func = OpFunction %void None %void_f\n"
+         "%label = OpLabel\n"
+         "OpReturn\n"
+         "OpFunctionEnd\n",
+         // expected
+         "OpCapability Shader\n"
+         "OpCapability VulkanMemoryModelKHR\n"
+         "OpExtension \"SPV_KHR_vulkan_memory_model\"\n"
+         "OpMemoryModel Logical VulkanKHR\n"
+         "OpEntryPoint Vertex %1 \"shader\"\n"
+         "%void = OpTypeVoid\n"
+         "%5 = OpTypeFunction %void\n"
+         "%1 = OpFunction %void None %5\n"
+         "%6 = OpLabel\n"
+         "OpReturn\n"
+         "OpFunctionEnd\n",
+         // pass
+         "strip-debug"}}));
 
 }  // namespace
 }  // namespace opt

3rdparty/spirv-tools/test/reduce/merge_blocks_test.cpp

@@ -508,6 +508,144 @@ TEST(MergeBlocksReductionPassTest, MergeWithOpPhi) {
   CheckEqual(env, after, context.get());
 }
 
+void MergeBlocksReductionPassTest_LoopReturn_Helper(bool reverse) {
+  // A merge block opportunity stores a block that can be merged with its
+  // predecessor.
+  // Given blocks A -> B -> C:
+  // This test demonstrates how merging B->C can invalidate
+  // the opportunity of merging A->B, and vice-versa. E.g.
+  // B->C are merged: B is now terminated with OpReturn.
+  // A->B can now no longer be merged because A is a loop header, which
+  // cannot be terminated with OpReturn.
+
+  std::string shader = R"(
+               OpCapability Shader
+          %1 = OpExtInstImport "GLSL.std.450"
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Fragment %2 "main"
+               OpExecutionMode %2 OriginUpperLeft
+               OpSource ESSL 310
+               OpName %2 "main"
+          %3 = OpTypeVoid
+          %4 = OpTypeFunction %3
+          %5 = OpTypeInt 32 1
+          %6 = OpTypePointer Function %5
+          %7 = OpTypeBool
+          %8 = OpConstantFalse %7
+          %2 = OpFunction %3 None %4
+          %9 = OpLabel
+               OpBranch %10
+         %10 = OpLabel                   ; A (loop header)
+               OpLoopMerge %13 %12 None
+               OpBranch %11
+         %12 = OpLabel                   ; (unreachable continue block)
+               OpBranch %10
+         %11 = OpLabel                   ; B
+               OpBranch %15
+         %15 = OpLabel                   ; C
+               OpReturn
+         %13 = OpLabel                   ; (unreachable merge block)
+               OpReturn
+               OpFunctionEnd
+  )";
+  const auto env = SPV_ENV_UNIVERSAL_1_3;
+  const auto consumer = nullptr;
+  const auto context =
+      BuildModule(env, consumer, shader, kReduceAssembleOption);
+  ASSERT_NE(context.get(), nullptr);
+  auto opportunities =
+      MergeBlocksReductionOpportunityFinder().GetAvailableOpportunities(
+          context.get());
+
+  // A->B and B->C
+  ASSERT_EQ(opportunities.size(), 2);
+
+  // Test applying opportunities in both orders.
+  if (reverse) {
+    std::reverse(opportunities.begin(), opportunities.end());
+  }
+
+  size_t num_applied = 0;
+  for (auto& ri : opportunities) {
+    if (ri->PreconditionHolds()) {
+      ri->TryToApply();
+      ++num_applied;
+    }
+  }
+
+  // Only 1 opportunity can be applied, as both disable each other.
+  ASSERT_EQ(num_applied, 1);
+
+  std::string after = R"(
+               OpCapability Shader
+          %1 = OpExtInstImport "GLSL.std.450"
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Fragment %2 "main"
+               OpExecutionMode %2 OriginUpperLeft
+               OpSource ESSL 310
+               OpName %2 "main"
+          %3 = OpTypeVoid
+          %4 = OpTypeFunction %3
+          %5 = OpTypeInt 32 1
+          %6 = OpTypePointer Function %5
+          %7 = OpTypeBool
+          %8 = OpConstantFalse %7
+          %2 = OpFunction %3 None %4
+          %9 = OpLabel
+               OpBranch %10
+         %10 = OpLabel                   ; A-B (loop header)
+               OpLoopMerge %13 %12 None
+               OpBranch %15
+         %12 = OpLabel                   ; (unreachable continue block)
+               OpBranch %10
+         %15 = OpLabel                   ; C
+               OpReturn
+         %13 = OpLabel                   ; (unreachable merge block)
+               OpReturn
+               OpFunctionEnd
+  )";
+
+  // The only difference is the labels.
+  std::string after_reversed = R"(
+               OpCapability Shader
+          %1 = OpExtInstImport "GLSL.std.450"
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Fragment %2 "main"
+               OpExecutionMode %2 OriginUpperLeft
+               OpSource ESSL 310
+               OpName %2 "main"
+          %3 = OpTypeVoid
+          %4 = OpTypeFunction %3
+          %5 = OpTypeInt 32 1
+          %6 = OpTypePointer Function %5
+          %7 = OpTypeBool
+          %8 = OpConstantFalse %7
+          %2 = OpFunction %3 None %4
+          %9 = OpLabel
+               OpBranch %10
+         %10 = OpLabel                   ; A (loop header)
+               OpLoopMerge %13 %12 None
+               OpBranch %11
+         %12 = OpLabel                   ; (unreachable continue block)
+               OpBranch %10
+         %11 = OpLabel                   ; B-C
+               OpReturn
+         %13 = OpLabel                   ; (unreachable merge block)
+               OpReturn
+               OpFunctionEnd
+  )";
+
+  CheckEqual(env, reverse ? after_reversed : after, context.get());
+}
+
+TEST(MergeBlocksReductionPassTest, LoopReturn) {
+  MergeBlocksReductionPassTest_LoopReturn_Helper(false);
+}
+
+TEST(MergeBlocksReductionPassTest, LoopReturnReverse) {
+  MergeBlocksReductionPassTest_LoopReturn_Helper(true);
+}
+
 }  // namespace
 }  // namespace reduce
 }  // namespace spvtools

tools/shaderc/shaderc_hlsl.cpp

@@ -107,7 +107,7 @@ namespace bgfx { namespace hlsl
 
 			if (g_verbose)
 			{
-				char filePath[MAX_PATH];
+				char filePath[bx::kMaxFilePath];
 				GetModuleFileNameA( (HMODULE)s_d3dcompilerdll, filePath, sizeof(filePath) );
 				BX_TRACE("Loaded %s compiler (%s).", compiler->fileName, filePath);
 			}

tools/texturev/texturev.cpp

@@ -23,8 +23,6 @@
 #include <imgui/imgui.h>
 #include <bgfx_utils.h>
 
-#include <dirent.h>
-
 #include <tinystl/allocator.h>
 #include <tinystl/vector.h>
 namespace stl = tinystl;
@@ -671,68 +669,76 @@ struct View
 
 	void updateFileList(const bx::FilePath& _filePath)
 	{
-		DIR* dir = opendir(_filePath.get() );
+		bx::DirectoryReader dr;
 
-		if (NULL == dir)
+		if (bx::open(&dr, _filePath) )
+		{
+			m_path = _filePath;
+		}
+		else if (bx::open(&dr, _filePath.getPath() ) )
 		{
 			m_path = _filePath.getPath();
-			dir = opendir(m_path.get() );
 		}
 		else
 		{
-			m_path = _filePath;
+			DBG("File path `%s` not found.", _filePath.get() );
+			return;
 		}
 
-		if (NULL != dir)
+		bx::Error err;
+
+		while (err.isOk() )
 		{
-			for (dirent* item = readdir(dir); NULL != item; item = readdir(dir) )
+			bx::FileInfo fi;
+			bx::read(&dr, fi, &err);
+
+			if (err.isOk()
+			&&  bx::FileType::File == fi.type)
 			{
-				if (0 == (item->d_type & DT_DIR) )
+				bx::StringView ext = fi.filePath.getExt();
+
+				if (!ext.isEmpty() )
 				{
-					const bx::StringView fileName(item->d_name);
-					bx::StringView ext = bx::strRFind(fileName, '.');
-					if (!ext.isEmpty() )
+					ext.set(ext.getPtr()+1, ext.getTerm() );
+
+					bool supported = false;
+					for (uint32_t ii = 0; ii < BX_COUNTOF(s_supportedExt); ++ii)
 					{
-						ext.set(ext.getPtr()+1, fileName.getTerm() );
-						bool supported = false;
-						for (uint32_t ii = 0; ii < BX_COUNTOF(s_supportedExt); ++ii)
+						if (0 == bx::strCmpI(ext, s_supportedExt[ii]) )
 						{
-							if (0 == bx::strCmpI(ext, s_supportedExt[ii]) )
-							{
-								supported = true;
-								break;
-							}
+							supported = true;
+							break;
 						}
+					}
 
-						if (supported)
-						{
-							m_fileList.push_back(item->d_name);
-						}
+					if (supported)
+					{
+						m_fileList.push_back(fi.filePath.get() );
 					}
 				}
 			}
+		}
+
+		bx::close(&dr);
 
-			std::sort(m_fileList.begin(), m_fileList.end(), sortNameAscending);
+		std::sort(m_fileList.begin(), m_fileList.end(), sortNameAscending);
 
-			m_fileIndex = 0;
-			uint32_t idx = 0;
-			for (FileList::const_iterator it = m_fileList.begin(); it != m_fileList.end(); ++it, ++idx)
+		m_fileIndex = 0;
+		uint32_t idx = 0;
+		for (FileList::const_iterator it = m_fileList.begin(); it != m_fileList.end(); ++it, ++idx)
+		{
+			if (0 == bx::strCmpI(it->c_str(), _filePath.getFileName() ) )
 			{
-				if (0 == bx::strCmpI(it->c_str(), _filePath.getFileName() ) )
-				{
-					// If it is case-insensitive match then might be correct one, but keep
-					// searching.
-					m_fileIndex = idx;
+				// If it is case-insensitive match then might be correct one, but keep
+				// searching.
+				m_fileIndex = idx;
 
-					if (0 == bx::strCmp(it->c_str(), _filePath.getFileName() ) )
-					{
-						// If it is exact match we're done.
-						break;
-					}
+				if (0 == bx::strCmp(it->c_str(), _filePath.getFileName() ) )
+				{
+					// If it is exact match we're done.
+					break;
 				}
 			}
-
-			closedir(dir);
 		}
 	}
 
@@ -1085,13 +1091,18 @@ void keyBindingHelp(const char* _bindings, const char* _description)
 	ImGui::Text(_description);
 }
 
+#if BX_PLATFORM_WINDOWS
+extern "C" void*    _stdcall GetModuleHandleA(const char* _moduleName);
+extern "C" uint32_t _stdcall GetModuleFileNameA(void* _module, char* _outFilePath, uint32_t _size);
+#endif // BX_PLATFORM_WINDOWS
+
 void associate()
 {
 #if BX_PLATFORM_WINDOWS
 	std::string str;
 
-	char exec[MAX_PATH];
-	GetModuleFileNameA(GetModuleHandleA(NULL), exec, MAX_PATH);
+	char exec[bx::kMaxFilePath];
+	GetModuleFileNameA(GetModuleHandleA(NULL), exec, sizeof(exec) );
 
 	std::string strExec = bx::replaceAll<std::string>(exec, "\\", "\\\\");