Add test for false loop init.

Clean up how for loop variables are declared.

reference/opt/shaders-msl/comp/torture-loop.comp

@@ -47,7 +47,8 @@ kernel void main0(uint3 gl_GlobalInvocationID [[thread_position_in_grid]], devic
     _96 = _40;
     float4 _100;
     uint _101;
-    for (uint _97 = 0u, _99 = _98; _97 < 16u; _95 = _100, _96++, _97++, _99 = _101)
+    uint _99;
+    for (uint _97 = 0u; _97 < 16u; _95 = _100, _96++, _97++, _99 = _101)
     {
         _100 = _95;
         _101 = 0u;

reference/opt/shaders-msl/frag/false-loop-init.frag

@@ -0,0 +1,38 @@
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+constant uint _49 = {};
+
+struct main0_in
+{
+    float4 accum [[user(locn0)]];
+};
+
+struct main0_out
+{
+    float4 result [[color(0)]];
+};
+
+fragment main0_out main0(main0_in in [[stage_in]])
+{
+    main0_out out = {};
+    out.result = float4(0.0);
+    uint _51;
+    uint _50;
+    for (int _48 = 0; _48 < 4; _48 += int(_51), _50 = _51)
+    {
+        if (in.accum.y > 10.0)
+        {
+            _51 = 40u;
+        }
+        else
+        {
+            _51 = 30u;
+        }
+        out.result += in.accum;
+    }
+    return out;
+}
+

reference/opt/shaders/comp/torture-loop.comp

@@ -45,7 +45,8 @@ void main()
     _96 = _40;
     vec4 _100;
     uint _101;
-    for (uint _97 = 0u, _99 = _98; _97 < 16u; _95 = _100, _96++, _97++, _99 = _101)
+    uint _99;
+    for (uint _97 = 0u; _97 < 16u; _95 = _100, _96++, _97++, _99 = _101)
     {
         _100 = _95;
         _101 = 0u;

reference/opt/shaders/flatten/multi-dimensional.desktop.flatten_dim.frag

@@ -13,7 +13,9 @@ void main()
     vec4 values3[2 * 3 * 1];
     int _96;
     int _97;
-    for (int _92 = 0, _94 = _93, _95 = _93; _92 < 2; _92++, _94 = _96, _95 = _97)
+    int _94;
+    int _95;
+    for (int _92 = 0; _92 < 2; _92++, _94 = _96, _95 = _97)
     {
         _96 = 0;
         _97 = _95;

reference/opt/shaders/frag/false-loop-init.frag

@@ -0,0 +1,28 @@
+#version 310 es
+precision mediump float;
+precision highp int;
+
+layout(location = 0) out vec4 result;
+layout(location = 0) in vec4 accum;
+
+uint _49;
+
+void main()
+{
+    result = vec4(0.0);
+    uint _51;
+    uint _50;
+    for (int _48 = 0; _48 < 4; _48 += int(_51), _50 = _51)
+    {
+        if (accum.y > 10.0)
+        {
+            _51 = 40u;
+        }
+        else
+        {
+            _51 = 30u;
+        }
+        result += accum;
+    }
+}
+

reference/shaders-msl/frag/false-loop-init.frag

@@ -0,0 +1,35 @@
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+struct main0_in
+{
+    float4 accum [[user(locn0)]];
+};
+
+struct main0_out
+{
+    float4 result [[color(0)]];
+};
+
+fragment main0_out main0(main0_in in [[stage_in]])
+{
+    main0_out out = {};
+    out.result = float4(0.0);
+    uint j;
+    for (int i = 0; i < 4; i += int(j))
+    {
+        if (in.accum.y > 10.0)
+        {
+            j = 40u;
+        }
+        else
+        {
+            j = 30u;
+        }
+        out.result += in.accum;
+    }
+    return out;
+}
+

reference/shaders/frag/false-loop-init.frag

@@ -0,0 +1,25 @@
+#version 310 es
+precision mediump float;
+precision highp int;
+
+layout(location = 0) out vec4 result;
+layout(location = 0) in vec4 accum;
+
+void main()
+{
+    result = vec4(0.0);
+    mediump uint j;
+    for (mediump int i = 0; i < 4; i += int(j))
+    {
+        if (accum.y > 10.0)
+        {
+            j = 40u;
+        }
+        else
+        {
+            j = 30u;
+        }
+        result += accum;
+    }
+}
+

shaders-msl/frag/false-loop-init.frag

@@ -0,0 +1,19 @@
+#version 310 es
+precision mediump float;
+
+layout(location = 0) in vec4 accum;
+layout(location = 0) out vec4 result;
+
+void main()
+{
+	result = vec4(0.0);
+	uint j;
+	for (int i = 0; i < 4; i += int(j))
+	{
+		if (accum.y > 10.0)
+			j = 40u;
+		else
+			j = 30u;
+		result += accum;
+	}
+}

shaders/frag/false-loop-init.frag

@@ -0,0 +1,19 @@
+#version 310 es
+precision mediump float;
+
+layout(location = 0) in vec4 accum;
+layout(location = 0) out vec4 result;
+
+void main()
+{
+	result = vec4(0.0);
+	uint j;
+	for (int i = 0; i < 4; i += int(j))
+	{
+		if (accum.y > 10.0)
+			j = 40u;
+		else
+			j = 30u;
+		result += accum;
+	}
+}

spirv_glsl.cpp

@@ -7181,9 +7181,17 @@ string CompilerGLSL::variable_decl(const SPIRVariable &variable)
 	auto res = join(to_qualifiers_glsl(variable.self), variable_decl(type, to_name(variable.self), variable.self));
 
 	if (variable.loop_variable && variable.static_expression)
-		res += join(" = ", to_expression(variable.static_expression));
+	{
+		uint32_t expr = variable.static_expression;
+		if (ids[expr].get_type() != TypeUndef)
+			res += join(" = ", to_expression(variable.static_expression));
+	}
 	else if (variable.initializer)
-		res += join(" = ", to_expression(variable.initializer));
+	{
+		uint32_t expr = variable.initializer;
+		if (ids[expr].get_type() != TypeUndef)
+			res += join(" = ", to_expression(variable.initializer));
+	}
 	return res;
 }
 
@@ -8023,16 +8031,22 @@ string CompilerGLSL::emit_for_loop_initializers(const SPIRBlock &block)
 	// If we cannot do this, declare individual variables before the loop header.
 
 	// We might have a loop variable candidate which was not assigned to for some reason.
-	bool missing_initializer = false;
+	uint32_t missing_initializers = 0;
 	for (auto &variable : block.loop_variables)
-		if (get<SPIRVariable>(variable).static_expression == 0)
-			missing_initializer = true;
+	{
+		uint32_t expr = get<SPIRVariable>(variable).static_expression;
 
-	if (block.loop_variables.size() == 1 && !missing_initializer)
+		// Sometimes loop variables are initialized with OpUndef, but we can just declare
+		// a plain variable without initializer in this case.
+		if (expr == 0 || ids[expr].get_type() == TypeUndef)
+			missing_initializers++;
+	}
+
+	if (block.loop_variables.size() == 1 && missing_initializers == 0)
 	{
 		return variable_decl(get<SPIRVariable>(block.loop_variables.front()));
 	}
-	else if (!same_types || missing_initializer)
+	else if (!same_types || missing_initializers == uint32_t(block.loop_variables.size()))
 	{
 		for (auto &loop_var : block.loop_variables)
 			statement(variable_decl(get<SPIRVariable>(loop_var)), ";");
@@ -8040,20 +8054,32 @@ string CompilerGLSL::emit_for_loop_initializers(const SPIRBlock &block)
 	}
 	else
 	{
-		auto &var = get<SPIRVariable>(block.loop_variables.front());
-		auto &type = get<SPIRType>(var.basetype);
-
-		// Don't remap the type here as we have multiple names,
-		// doesn't make sense to remap types for loop variables anyways.
-		// It is assumed here that all relevant qualifiers are equal for all loop variables.
-		string expr = join(to_qualifiers_glsl(var.self), type_to_glsl(type), " ");
+		// We have a mix of loop variables, either ones with a clear initializer, or ones without.
+		// Separate the two streams.
+		string expr;
 
 		for (auto &loop_var : block.loop_variables)
 		{
-			auto &v = get<SPIRVariable>(loop_var);
-			expr += join(to_name(loop_var), " = ", to_expression(v.static_expression));
-			if (&loop_var != &block.loop_variables.back())
-				expr += ", ";
+			uint32_t static_expr = get<SPIRVariable>(loop_var).static_expression;
+			if (static_expr == 0 || ids[static_expr].get_type() == TypeUndef)
+			{
+				statement(variable_decl(get<SPIRVariable>(loop_var)), ";");
+			}
+			else
+			{
+				if (expr.empty())
+				{
+					// For loop initializers are of the form <type id = value, id = value, id = value, etc ...
+					auto &var = get<SPIRVariable>(block.loop_variables.front());
+					auto &type = get<SPIRType>(var.basetype);
+					expr = join(to_qualifiers_glsl(var.self), type_to_glsl(type), " ");
+				}
+				else
+					expr += ", ";
+
+				auto &v = get<SPIRVariable>(loop_var);
+				expr += join(to_name(loop_var), " = ", to_expression(v.static_expression));
+			}
 		}
 		return expr;
 	}
@@ -8064,11 +8090,21 @@ bool CompilerGLSL::for_loop_initializers_are_same_type(const SPIRBlock &block)
 	if (block.loop_variables.size() <= 1)
 		return true;
 
-	uint32_t expected = get<SPIRVariable>(block.loop_variables[0]).basetype;
-	uint64_t expected_flags = get_decoration_mask(block.loop_variables[0]);
+	uint32_t expected = 0;
+	uint64_t expected_flags = 0;
 	for (auto &var : block.loop_variables)
 	{
-		if (expected != get<SPIRVariable>(var).basetype)
+		// Don't care about uninitialized variables as they will not be part of the initializers.
+		uint32_t expr = get<SPIRVariable>(var).static_expression;
+		if (expr == 0 || ids[expr].get_type() == TypeUndef)
+			continue;
+
+		if (expected == 0)
+		{
+			expected = get<SPIRVariable>(var).basetype;
+			expected_flags = get_decoration_mask(var);
+		}
+		else if (expected != get<SPIRVariable>(var).basetype)
 			return false;
 
 		// Precision flags and things like that must also match.