Merge branch 'msl_global-invocation-id' of git://github.com/Kangz/SPIRV-Cross into pr-105

reference/shaders-msl/comp/global-invocation-id-writable-ssbo-in-function.comp

@@ -0,0 +1,31 @@
+#pragma clang diagnostic ignored "-Wmissing-prototypes"
+
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+struct myBlock
+{
+    int a;
+    float b[1];
+};
+
+// Support GLSL mod(), which is slightly different than Metal fmod()
+template<typename Tx, typename Ty>
+Tx mod(Tx x, Ty y)
+{
+    return x - y * floor(x / y);
+}
+
+float getB(device const myBlock& myStorage, thread const uint3& gl_GlobalInvocationID)
+{
+    return myStorage.b[gl_GlobalInvocationID.x];
+}
+
+kernel void main0(device myBlock& myStorage [[buffer(0)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
+{
+    myStorage.a = (myStorage.a + 1) % 256;
+    myStorage.b[gl_GlobalInvocationID.x] = mod(getB(myStorage, gl_GlobalInvocationID) + 0.0199999995529651641845703125, 1.0);
+}
+

reference/shaders-msl/comp/global-invocation-id.comp

@@ -0,0 +1,24 @@
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+struct myBlock
+{
+    int a;
+    float b[1];
+};
+
+// Support GLSL mod(), which is slightly different than Metal fmod()
+template<typename Tx, typename Ty>
+Tx mod(Tx x, Ty y)
+{
+    return x - y * floor(x / y);
+}
+
+kernel void main0(device myBlock& myStorage [[buffer(0)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
+{
+    myStorage.a = (myStorage.a + 1) % 256;
+    myStorage.b[gl_GlobalInvocationID.x] = mod(myStorage.b[gl_GlobalInvocationID.x] + 0.0199999995529651641845703125, 1.0);
+}
+

reference/shaders-msl/comp/local-invocation-id.comp

@@ -0,0 +1,24 @@
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+struct myBlock
+{
+    int a;
+    float b[1];
+};
+
+// Support GLSL mod(), which is slightly different than Metal fmod()
+template<typename Tx, typename Ty>
+Tx mod(Tx x, Ty y)
+{
+    return x - y * floor(x / y);
+}
+
+kernel void main0(device myBlock& myStorage [[buffer(0)]], uint3 gl_LocalInvocationID [[thread_position_in_threadgroup]])
+{
+    myStorage.a = (myStorage.a + 1) % 256;
+    myStorage.b[gl_LocalInvocationID.x] = mod(myStorage.b[gl_LocalInvocationID.x] + 0.0199999995529651641845703125, 1.0);
+}
+

reference/shaders-msl/comp/local-invocation-index.comp

@@ -0,0 +1,24 @@
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+struct myBlock
+{
+    int a;
+    float b[1];
+};
+
+// Support GLSL mod(), which is slightly different than Metal fmod()
+template<typename Tx, typename Ty>
+Tx mod(Tx x, Ty y)
+{
+    return x - y * floor(x / y);
+}
+
+kernel void main0(device myBlock& myStorage [[buffer(0)]], uint gl_LocalInvocationIndex [[thread_index_in_threadgroup]])
+{
+    myStorage.a = (myStorage.a + 1) % 256;
+    myStorage.b[gl_LocalInvocationIndex] = mod(myStorage.b[gl_LocalInvocationIndex] + 0.0199999995529651641845703125, 1.0);
+}
+

reference/shaders-msl/comp/writable-ssbo.comp

@@ -0,0 +1,24 @@
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+struct myBlock
+{
+    int a;
+    float b;
+};
+
+// Support GLSL mod(), which is slightly different than Metal fmod()
+template<typename Tx, typename Ty>
+Tx mod(Tx x, Ty y)
+{
+    return x - y * floor(x / y);
+}
+
+kernel void main0(device myBlock& myStorage [[buffer(0)]])
+{
+    myStorage.a = (myStorage.a + 1) % 256;
+    myStorage.b = mod(myStorage.b + 0.0199999995529651641845703125, 1.0);
+}
+

shaders-msl/comp/global-invocation-id-writable-ssbo-in-function.comp

@@ -0,0 +1,12 @@
+#version 450
+layout(set = 0, binding = 0) buffer myBlock {
+    int a;
+    float b[1];
+} myStorage;
+float getB() {
+    return myStorage.b[gl_GlobalInvocationID.x];
+}
+void main() {
+    myStorage.a = (myStorage.a + 1) % 256;
+    myStorage.b[gl_GlobalInvocationID.x] = mod((getB() + 0.02), 1.0);
+}

shaders-msl/comp/global-invocation-id.comp

@@ -0,0 +1,9 @@
+#version 450
+layout(set = 0, binding = 0) buffer myBlock {
+    int a;
+    float b[1];
+} myStorage;
+void main() {
+    myStorage.a = (myStorage.a + 1) % 256;
+    myStorage.b[gl_GlobalInvocationID.x] = mod((myStorage.b[gl_GlobalInvocationID.x] + 0.02), 1.0);
+}

shaders-msl/comp/local-invocation-id.comp

@@ -0,0 +1,9 @@
+#version 450
+layout(set = 0, binding = 0) buffer myBlock {
+    int a;
+    float b[1];
+} myStorage;
+void main() {
+    myStorage.a = (myStorage.a + 1) % 256;
+    myStorage.b[gl_LocalInvocationID.x] = mod((myStorage.b[gl_LocalInvocationID.x] + 0.02), 1.0);
+}

shaders-msl/comp/local-invocation-index.comp

@@ -0,0 +1,9 @@
+#version 450
+layout(set = 0, binding = 0) buffer myBlock {
+    int a;
+    float b[1];
+} myStorage;
+void main() {
+    myStorage.a = (myStorage.a + 1) % 256;
+    myStorage.b[gl_LocalInvocationIndex.x] = mod((myStorage.b[gl_LocalInvocationIndex.x] + 0.02), 1.0);
+}

shaders-msl/comp/writable-ssbo.comp

@@ -0,0 +1,9 @@
+#version 450
+layout(set = 0, binding = 0) buffer myBlock {
+    int a;
+    float b;
+} myStorage;
+void main() {
+    myStorage.a = (myStorage.a + 1) % 256;
+    myStorage.b = mod((myStorage.b + 0.02), 1.0);
+}

spirv_msl.cpp

@@ -16,7 +16,9 @@
 
 #include "spirv_msl.hpp"
 #include "GLSL.std.450.h"
+
 #include <algorithm>
+#include <cassert>
 #include <numeric>
 
 using namespace spv;
@@ -717,22 +719,16 @@ void CompilerMSL::emit_function_prototype(SPIRFunction &func, uint64_t)
 	{
 		add_local_variable_name(arg.id);
 
-		bool is_uniform_struct = false;
+		string address_space = "thread";
+
 		auto *var = maybe_get<SPIRVariable>(arg.id);
 		if (var)
 		{
 			var->parameter = &arg; // Hold a pointer to the parameter so we can invalidate the readonly field if needed.
-
-			// Check if this arg is one of the synthetic uniform args
-			// created to handle uniform access inside the function
-			auto &var_type = get<SPIRType>(var->basetype);
-			is_uniform_struct =
-			    ((var_type.basetype == SPIRType::Struct) &&
-			     (var_type.storage == StorageClassUniform || var_type.storage == StorageClassUniformConstant ||
-			      var_type.storage == StorageClassPushConstant));
+			address_space = get_argument_address_space(*var);
 		}
 
-		decl += (is_uniform_struct ? "constant " : "thread ");
+		decl += address_space + " ";
 		decl += argument_decl(arg);
 
 		// Manufacture automatic sampler arg for SampledImage texture
@@ -1154,6 +1150,24 @@ string CompilerMSL::member_attribute_qualifier(const SPIRType &type, uint32_t in
 		return string(" [[color(") + convert_to_string(locn) + ")]]";
 	}
 
+	// Compute function inputs
+	if (execution.model == ExecutionModelGLCompute && type.storage == StorageClassInput)
+	{
+		if (is_builtin)
+		{
+			switch (builtin)
+			{
+			case BuiltInGlobalInvocationId:
+			case BuiltInLocalInvocationId:
+			case BuiltInLocalInvocationIndex:
+				return string(" [[") + builtin_qualifier(builtin) + "]]";
+
+			default:
+				return "";
+			}
+		}
+	}
+
 	return "";
 }
 
@@ -1256,6 +1270,29 @@ string CompilerMSL::clean_func_name(string func_name)
 	return (iter != func_name_overrides.end()) ? iter->second : func_name;
 }
 
+// In MSL address space qualifiers are required for all pointer or reference arguments
+string CompilerMSL::get_argument_address_space(const SPIRVariable &argument)
+{
+	const auto &type = get<SPIRType>(argument.basetype);
+
+	if ((type.basetype == SPIRType::Struct) &&
+	    (type.storage == StorageClassUniform || type.storage == StorageClassUniformConstant ||
+	     type.storage == StorageClassPushConstant))
+	{
+		if ((meta[type.self].decoration.decoration_flags & (1ull << DecorationBufferBlock)) != 0 &&
+		    (meta[argument.self].decoration.decoration_flags & (1ull << DecorationNonWritable)) == 0)
+		{
+			return "device";
+		}
+		else
+		{
+			return "constant";
+		}
+	}
+
+	return "thread";
+}
+
 // Returns a string containing a comma-delimited list of args for the entry point function
 string CompilerMSL::entry_point_args(bool append_comma)
 {
@@ -1293,16 +1330,7 @@ string CompilerMSL::entry_point_args(bool append_comma)
 						break;
 					if (!ep_args.empty())
 						ep_args += ", ";
-					if ((meta[type.self].decoration.decoration_flags & (1ull << DecorationBufferBlock)) != 0 &&
-					    (meta[var.self].decoration.decoration_flags & (1ull << DecorationNonWritable)) == 0)
-					{
-						ep_args += "device ";
-					}
-					else
-					{
-						ep_args += "constant ";
-					}
-					ep_args += type_to_glsl(type) + "& " + to_name(var.self);
+					ep_args += get_argument_address_space(var) + " " + type_to_glsl(type) + "& " + to_name(var.self);
 					ep_args += " [[buffer(" + convert_to_string(get_metal_resource_index(var, type.basetype)) + ")]]";
 					break;
 				}
@@ -1673,6 +1701,16 @@ string CompilerMSL::builtin_qualifier(BuiltIn builtin)
 			return "depth(any)";
 	}
 
+	// Compute function in
+	case BuiltInGlobalInvocationId:
+		return "thread_position_in_grid";
+
+	case BuiltInLocalInvocationId:
+		return "thread_position_in_threadgroup";
+
+	case BuiltInLocalInvocationIndex:
+		return "thread_index_in_threadgroup";
+
 	default:
 		return "unsupported-built-in";
 	}
@@ -1713,6 +1751,14 @@ string CompilerMSL::builtin_type_decl(BuiltIn builtin)
 	case BuiltInSampleMask:
 		return "uint";
 
+	// Compute function in
+	case BuiltInGlobalInvocationId:
+		return "uint3";
+	case BuiltInLocalInvocationId:
+		return "uint3";
+	case BuiltInLocalInvocationIndex:
+		return "uint";
+
 	default:
 		return "unsupported-built-in-type";
 	}

spirv_msl.hpp

@@ -117,6 +117,8 @@ protected:
 	                             uint32_t grad_y, uint32_t lod, uint32_t coffset, uint32_t offset, uint32_t bias,
 	                             uint32_t comp, uint32_t sample, bool *p_forward) override;
 
+	std::string get_argument_address_space(const SPIRVariable &argument);
+
 	void preprocess_op_codes();
 	void emit_custom_functions();
 	void localize_global_variables();