Deal with scoping for Private variables.

reference/opt/shaders-msl/asm/comp/storage-buffer-basic.invalid.asm.comp

@@ -16,7 +16,7 @@ struct _6
 
 kernel void main0(device _6& _8 [[buffer(0)]], device _6& _9 [[buffer(1)]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]])
 {
-    uint3 _23 = gl_WorkGroupSize;
     _8._m0[gl_WorkGroupID.x] = _9._m0[gl_WorkGroupID.x] + _8._m0[gl_WorkGroupID.x];
+    uint3 _23 = gl_WorkGroupSize;
 }
 

reference/opt/shaders-msl/comp/access-private-workgroup-in-function.comp

@@ -0,0 +1,18 @@
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+kernel void main0(uint gl_LocalInvocationIndex [[thread_index_in_threadgroup]])
+{
+    threadgroup int u;
+    u = 50;
+    if (gl_LocalInvocationIndex == 0u)
+    {
+    }
+    else
+    {
+        u = 20;
+    }
+}
+

reference/shaders-msl/asm/comp/storage-buffer-basic.invalid.asm.comp

@@ -16,7 +16,7 @@ struct _6
 
 kernel void main0(device _6& _8 [[buffer(0)]], device _6& _9 [[buffer(1)]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]])
 {
-    uint3 _23 = gl_WorkGroupSize;
     _8._m0[gl_WorkGroupID.x] = _9._m0[gl_WorkGroupID.x] + _8._m0[gl_WorkGroupID.x];
+    uint3 _23 = gl_WorkGroupSize;
 }
 

reference/shaders-msl/comp/access-private-workgroup-in-function.comp

@@ -0,0 +1,34 @@
+#pragma clang diagnostic ignored "-Wmissing-prototypes"
+
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+void set_f(thread int& f)
+{
+    f = 40;
+}
+
+void set_shared_u(threadgroup int& u)
+{
+    u = 50;
+}
+
+kernel void main0(uint gl_LocalInvocationIndex [[thread_index_in_threadgroup]])
+{
+    threadgroup int u;
+    int f;
+    set_f(f);
+    set_shared_u(u);
+    if (gl_LocalInvocationIndex == 0u)
+    {
+        f = 10;
+    }
+    else
+    {
+        f = 30;
+        u = 20;
+    }
+}
+

reference/shaders-msl/vert/packed_matrix.vert

@@ -40,6 +40,7 @@ struct main0_out
 vertex main0_out main0(main0_in in [[stage_in]], constant _42& _44 [[buffer(12)]], constant _15& _17 [[buffer(13)]])
 {
     main0_out out = {};
+    float3 _91;
     float3 _13;
     do
     {

shaders-msl/comp/access-private-workgroup-in-function.comp

@@ -0,0 +1,31 @@
+#version 450
+layout(local_size_x = 1) in;
+
+int f;
+shared int u;
+
+void set_f()
+{
+	f = 40;
+}
+
+void set_shared_u()
+{
+	u = 50;
+}
+
+void main()
+{
+	set_f();
+	set_shared_u();
+	if (gl_LocalInvocationIndex == 0u)
+	{
+		f = 10;
+	}
+	else
+	{
+		f = 30;
+		u = 20;
+	}
+}
+

spirv_glsl.cpp

@@ -8862,6 +8862,20 @@ void CompilerGLSL::emit_function(SPIRFunction &func, const Bitset &return_flags)
 			statement(variable_decl(var), ";");
 			var.deferred_declaration = false;
 		}
+		else if (var.storage == StorageClassPrivate)
+		{
+			// These variables will not have had their CFG usage analyzed, so move it to the entry block.
+			// Comes from MSL which can push global variables as local variables in main function.
+			// We could just declare them right now, but we would miss out on an important initialization case which is
+			// LUT declaration in MSL.
+			// If we don't declare the variable when it is assigned we're forced to go through a helper function
+			// which copies elements one by one.
+			add_local_variable_name(var.self);
+			auto &dominated = entry_block.dominated_variables;
+			if (find(begin(dominated), end(dominated), var.self) == end(dominated))
+				entry_block.dominated_variables.push_back(var.self);
+			var.deferred_declaration = true;
+		}
 		else if (expression_is_lvalue(v))
 		{
 			add_local_variable_name(var.self);

spirv_msl.cpp

@@ -416,7 +416,6 @@ void CompilerMSL::resolve_specialized_array_lengths()
 // extract that variable and add it as an argument to that function.
 void CompilerMSL::extract_global_variables_from_functions()
 {
-
 	// Uniforms
 	unordered_set<uint32_t> global_var_ids;
 	for (auto &id : ids)
@@ -433,10 +432,11 @@ void CompilerMSL::extract_global_variables_from_functions()
 		}
 	}
 
-	// Local vars that are declared in the main function and accessed directy by a function
+	// Local vars that are declared in the main function and accessed directly by a function
 	auto &entry_func = get<SPIRFunction>(entry_point);
 	for (auto &var : entry_func.local_variables)
-		global_var_ids.insert(var);
+		if (get<SPIRVariable>(var).storage != StorageClassFunction)
+			global_var_ids.insert(var);
 
 	std::set<uint32_t> added_arg_ids;
 	unordered_set<uint32_t> processed_func_ids;
@@ -491,6 +491,7 @@ void CompilerMSL::extract_global_variables_from_function(uint32_t func_id, std::
 
 				break;
 			}
+
 			case OpFunctionCall:
 			{
 				// First see if any of the function call args are globals
@@ -510,9 +511,21 @@ void CompilerMSL::extract_global_variables_from_function(uint32_t func_id, std::
 				break;
 			}
 
+			case OpStore:
+			{
+				uint32_t base_id = ops[0];
+				if (global_var_ids.find(base_id) != global_var_ids.end())
+					added_arg_ids.insert(base_id);
+				break;
+			}
+
 			default:
 				break;
 			}
+
+			// TODO: Add all other operations which can affect memory.
+			// We should consider a more unified system here to reduce boiler-plate.
+			// This kind of analysis is done in several places ...
 		}
 	}