From 3c285a15e0a2747b4469627cc8a5bfdf23490a63 Mon Sep 17 00:00:00 2001
From: Hans-Kristian Arntzen <hans-kristian.arntzen@arm.com>
Date: Mon, 4 Jul 2016 13:30:05 +0200
Subject: [PATCH] Add support for reflecting execution modes.

---
 main.cpp        | 63 ++++++++++++++++++++++++++++++++++++++++++++++
 spirv_cross.cpp | 67 +++++++++++++++++++++++++++++++++++++++++++++++++
 spirv_cross.hpp | 11 ++++++++
 3 files changed, 141 insertions(+)

diff --git a/main.cpp b/main.cpp
index e79bc67c..bffa229d 100644
--- a/main.cpp
+++ b/main.cpp
@@ -212,6 +212,69 @@ static void print_resources(const Compiler &compiler, const char *tag, const vec
 
 static void print_resources(const Compiler &compiler, const ShaderResources &res)
 {
+	uint64_t modes = compiler.get_execution_mode_mask();
+
+	fprintf(stderr, "Execution modes:\n");
+	for (unsigned i = 0; i < 64; i++)
+	{
+		if (!(modes & (1ull << i)))
+			continue;
+
+		auto mode = static_cast<ExecutionMode>(i);
+		uint32_t arg0 = compiler.get_execution_mode_argument(mode, 0);
+		uint32_t arg1 = compiler.get_execution_mode_argument(mode, 1);
+		uint32_t arg2 = compiler.get_execution_mode_argument(mode, 2);
+
+		switch (static_cast<ExecutionMode>(i))
+		{
+			case ExecutionModeInvocations:
+				fprintf(stderr, "  Invocations: %u\n", arg0);
+				break;
+
+			case ExecutionModeLocalSize:
+				fprintf(stderr, "  LocalSize: (%u, %u, %u)\n", arg0, arg1, arg2);
+				break;
+
+			case ExecutionModeOutputVertices:
+				fprintf(stderr, "  OutputVertices: %u\n", arg0);
+				break;
+
+#define CHECK_MODE(m) case ExecutionMode##m: fprintf(stderr, "  %s\n", #m); break
+			    CHECK_MODE(SpacingEqual);
+			    CHECK_MODE(SpacingFractionalEven);
+			    CHECK_MODE(SpacingFractionalOdd);
+			    CHECK_MODE(VertexOrderCw);
+			    CHECK_MODE(VertexOrderCcw);
+			    CHECK_MODE(PixelCenterInteger);
+			    CHECK_MODE(OriginUpperLeft);
+			    CHECK_MODE(OriginLowerLeft);
+			    CHECK_MODE(EarlyFragmentTests);
+			    CHECK_MODE(PointMode);
+			    CHECK_MODE(Xfb);
+			    CHECK_MODE(DepthReplacing);
+			    CHECK_MODE(DepthGreater);
+			    CHECK_MODE(DepthLess);
+			    CHECK_MODE(DepthUnchanged);
+			    CHECK_MODE(LocalSizeHint);
+			    CHECK_MODE(InputPoints);
+			    CHECK_MODE(InputLines);
+			    CHECK_MODE(InputLinesAdjacency);
+			    CHECK_MODE(Triangles);
+			    CHECK_MODE(InputTrianglesAdjacency);
+			    CHECK_MODE(Quads);
+			    CHECK_MODE(Isolines);
+			    CHECK_MODE(OutputPoints);
+			    CHECK_MODE(OutputLineStrip);
+			    CHECK_MODE(OutputTriangleStrip);
+			    CHECK_MODE(VecTypeHint);
+			    CHECK_MODE(ContractionOff);
+
+			default:
+				break;
+		}
+	}
+	fprintf(stderr, "\n");
+
     print_resources(compiler, "subpass inputs", res.subpass_inputs);
     print_resources(compiler, "inputs", res.stage_inputs);
     print_resources(compiler, "outputs", res.stage_outputs);
diff --git a/spirv_cross.cpp b/spirv_cross.cpp
index 5a320cb4..2cd01f9c 100644
--- a/spirv_cross.cpp
+++ b/spirv_cross.cpp
@@ -1931,3 +1931,70 @@ bool Compiler::types_are_logically_equivalent(const SPIRType &a, const SPIRType
 
 	return true;
 }
+
+uint64_t Compiler::get_execution_mode_mask() const
+{
+	return execution.flags;
+}
+
+void Compiler::set_execution_mode(ExecutionMode mode, uint32_t arg0, uint32_t arg1, uint32_t arg2)
+{
+	execution.flags |= 1ull << mode;
+	switch (mode)
+	{
+	case ExecutionModeLocalSize:
+		execution.workgroup_size.x = arg0;
+		execution.workgroup_size.y = arg1;
+		execution.workgroup_size.z = arg2;
+		break;
+
+	case ExecutionModeInvocations:
+		execution.invocations = arg0;
+		break;
+
+	case ExecutionModeOutputVertices:
+		execution.output_vertices = arg0;
+		break;
+
+	default:
+		break;
+	}
+}
+
+void Compiler::unset_execution_mode(ExecutionMode mode)
+{
+	execution.flags &= ~(1ull << mode);
+}
+
+uint32_t Compiler::get_execution_mode_argument(spv::ExecutionMode mode, uint32_t index) const
+{
+	switch (mode)
+	{
+	case ExecutionModeLocalSize:
+		switch (index)
+		{
+		case 0:
+			return execution.workgroup_size.x;
+		case 1:
+			return execution.workgroup_size.y;
+		case 2:
+			return execution.workgroup_size.z;
+		default:
+			return 0;
+		}
+
+	case ExecutionModeInvocations:
+		return execution.invocations;
+
+	case ExecutionModeOutputVertices:
+		return execution.output_vertices;
+
+	default:
+		return 0;
+	}
+}
+
+ExecutionModel Compiler::get_execution_model() const
+{
+	return execution.model;
+}
diff --git a/spirv_cross.hpp b/spirv_cross.hpp
index 28b04abc..c83d98f9 100644
--- a/spirv_cross.hpp
+++ b/spirv_cross.hpp
@@ -171,6 +171,17 @@ public:
 	// Query shader resources, use ids with reflection interface to modify or query binding points, etc.
 	ShaderResources get_shader_resources() const;
 
+	// Query and modify OpExecutionMode.
+	uint64_t get_execution_mode_mask() const;
+	void unset_execution_mode(spv::ExecutionMode mode);
+	void set_execution_mode(spv::ExecutionMode mode, uint32_t arg0 = 0, uint32_t arg1 = 0, uint32_t arg2 = 0);
+
+	// Gets argument for an execution mode (LocalSize, Invocations, OutputVertices).
+	// For LocalSize, the index argument is used to select the dimension (X = 0, Y = 1, Z = 2).
+	// For execution modes which do not have arguments, 0 is returned.
+	uint32_t get_execution_mode_argument(spv::ExecutionMode mode, uint32_t index = 0) const;
+	spv::ExecutionModel get_execution_model() const;
+
 protected:
 	const uint32_t *stream(const Instruction &instr) const
 	{