diff --git a/reference/opt/shaders-msl/asm/comp/storage-buffer-basic.invalid.asm.comp b/reference/opt/shaders-msl/asm/comp/storage-buffer-basic.invalid.asm.comp index 9e37362d..2c9b038b 100644 --- a/reference/opt/shaders-msl/asm/comp/storage-buffer-basic.invalid.asm.comp +++ b/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; } diff --git a/reference/opt/shaders-msl/comp/access-private-workgroup-in-function.comp b/reference/opt/shaders-msl/comp/access-private-workgroup-in-function.comp new file mode 100644 index 00000000..44405126 --- /dev/null +++ b/reference/opt/shaders-msl/comp/access-private-workgroup-in-function.comp @@ -0,0 +1,18 @@ +#include +#include + +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; + } +} + diff --git a/reference/shaders-msl/asm/comp/storage-buffer-basic.invalid.asm.comp b/reference/shaders-msl/asm/comp/storage-buffer-basic.invalid.asm.comp index 9e37362d..2c9b038b 100644 --- a/reference/shaders-msl/asm/comp/storage-buffer-basic.invalid.asm.comp +++ b/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; } diff --git a/reference/shaders-msl/comp/access-private-workgroup-in-function.comp b/reference/shaders-msl/comp/access-private-workgroup-in-function.comp new file mode 100644 index 00000000..17acda96 --- /dev/null +++ b/reference/shaders-msl/comp/access-private-workgroup-in-function.comp @@ -0,0 +1,34 @@ +#pragma clang diagnostic ignored "-Wmissing-prototypes" + +#include +#include + +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; + } +} + diff --git a/reference/shaders-msl/vert/packed_matrix.vert b/reference/shaders-msl/vert/packed_matrix.vert index e3e2ff96..1f26a68c 100644 --- a/reference/shaders-msl/vert/packed_matrix.vert +++ b/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 { diff --git a/shaders-msl/comp/access-private-workgroup-in-function.comp b/shaders-msl/comp/access-private-workgroup-in-function.comp new file mode 100644 index 00000000..7cb1e6f1 --- /dev/null +++ b/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; + } +} + diff --git a/spirv_glsl.cpp b/spirv_glsl.cpp index 90f41e1d..56085bea 100644 --- a/spirv_glsl.cpp +++ b/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); diff --git a/spirv_msl.cpp b/spirv_msl.cpp index 0e44e322..8db8a943 100644 --- a/spirv_msl.cpp +++ b/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 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(entry_point); for (auto &var : entry_func.local_variables) - global_var_ids.insert(var); + if (get(var).storage != StorageClassFunction) + global_var_ids.insert(var); std::set added_arg_ids; unordered_set 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 ... } }