From a347cbcc8284301ad6989d486373f0d2ffb9657c Mon Sep 17 00:00:00 2001
From: Sylvain BERTRAND <sylvain.bertrand@legeek.net>
Date: Sat, 16 Nov 2019 17:48:29 +0000
Subject: [PATCH] braindead partial spir-v disassembler

---
 spirv/dis/README |    8 +
 spirv/dis/dis.c  | 2353 ++++++++++++++++++++++++++++++++++++++++++++++++++++++
 2 files changed, 2361 insertions(+)
 create mode 100644 spirv/dis/README
 create mode 100644 spirv/dis/dis.c

diff --git a/spirv/dis/README b/spirv/dis/README
new file mode 100644
index 0000000..cc6a282
--- /dev/null
+++ b/spirv/dis/README
@@ -0,0 +1,8 @@
+braindead partial spir-v disassembler.
+
+PROS:
+    - it is not brainf*cked hate incarnation, namely not c++
+    - it is __simple__ C
+
+CONS:
+    - it is not RISC-V assembly
diff --git a/spirv/dis/dis.c b/spirv/dis/dis.c
new file mode 100644
index 0000000..c6ee9ab
--- /dev/null
+++ b/spirv/dis/dis.c
@@ -0,0 +1,2353 @@
+#include <stdbool.h>
+#include <stdio.h>
+#include <stdint.h>
+#include <stdlib.h>
+#include <string.h>
+/*
+ * ABBREVIATIONS:
+ * blk(s) : BLocK(S)
+ * cap(s) : CAPability(ieS)
+ * decl : DECLaration
+ * def(s) | DEFinition(S)
+ * ext : EXTented
+ * extn(s) : EXTensioN(S)
+ * hdr : HeaDeR
+ * id(s) : IDentifier(S)
+ * img(s) : IMaGe(S)
+ * inst(s) : INSTruction(S)
+ * intf(s) : INTerFace(S)
+ * lang :LANGuage
+ * lit(s): LITeral(S)
+ * litstr : LITeral STRing
+ * n : couNt
+ * num : NUMber
+ * op(s) : OPcode(S)
+ * opd(s) : OPeranD)S)
+ * src : SouRCe
+ * str : STRing
+ * sz : SiZe
+ * w(s) : Word(S)
+ */
+#define u8 uint8_t
+#define u16 uint16_t
+#define u32 uint32_t
+#define constant enum
+#define loop for(;;)
+#define out(fmt,...) fprintf(stdout,fmt, ##__VA_ARGS__)
+#define out_depth(fmt,...) \
+{ \
+	u8 d; \
+	d = 0; \
+	loop { \
+		if (d == depth) \
+			break; \
+		fprintf(stdout, "\t"); \
+		++d; \
+	} \
+} \
+fprintf(stdout,fmt, ##__VA_ARGS__)
+
+constant {
+	spirv_op_undef = 1,
+	spirv_op_sourcecontinued = 2,
+	spirv_op_source = 3,
+	spirv_op_sourceextension = 4,
+	spirv_op_name = 5,
+	spirv_op_membername = 6,
+	spirv_op_string = 7,
+	spirv_op_line = 8,
+	spirv_op_extension = 10,
+	spirv_op_extinstimport = 11,
+	spirv_op_extinst = 12,
+	spirv_op_memorymodel = 14,
+	spirv_op_entrypoint = 15,
+	spirv_op_executionmode = 16,
+	spirv_op_capability = 17,
+	spirv_op_typevoid = 19,
+	spirv_op_typebool = 20,
+	spirv_op_typeint = 21,
+	spirv_op_typefloat = 22,
+	spirv_op_typevector = 23,
+	spirv_op_typematrix = 24,
+	spirv_op_typeimage = 25,
+	spirv_op_typesampler = 26,
+	spirv_op_typesampledimage = 27,
+	spirv_op_typearray = 28,
+	spirv_op_typeruntimearray = 29,
+	spirv_op_typestruct= 30,
+	spirv_op_typeopaque= 31,
+	spirv_op_typepointer= 32,
+	spirv_op_typefunction= 33,
+	spirv_op_typeevent = 34,
+	spirv_op_typedeviceevent = 35,
+	spirv_op_typereserveid = 36,
+	spirv_op_typequeue = 37,
+	spirv_op_typepipe = 38,
+	spirv_op_typeforwardpointer = 39,
+	spirv_op_constanttrue = 41,
+	spirv_op_constantfalse = 42,
+	spirv_op_constant = 43,
+	spirv_op_constantcomposite = 44,
+	spirv_op_constantsampler = 45,
+	spirv_op_constantnull = 46,
+	spirv_op_specconstanttrue = 48,
+	spirv_op_specconstantfalse = 49,
+	spirv_op_specconstant = 50,
+	spirv_op_specconstantcomposite = 51,
+	spirv_op_specconstantop = 52,
+	spirv_op_function = 54,
+	spirv_op_functionparameter = 55,
+	spirv_op_functionend = 56,
+	spirv_op_functioncall = 57,
+	spirv_op_variable = 59,
+	spirv_op_load = 61,
+	spirv_op_store = 62,
+	spirv_op_accesschain = 65,
+	spirv_op_decorate = 71,
+	spirv_op_memberdecorate = 72,
+	spirv_op_decorationgroup = 73,
+	spirv_op_groupdecorate = 74,
+	spirv_op_groupmemberdecorate = 75,
+	spirv_op_vectorshuffle = 79,
+	spirv_op_compositeconstruct = 80,
+	spirv_op_compositeextract = 81,
+	spirv_op_compositeinsert = 82,
+	spirv_op_sampledimage = 86,
+	spirv_op_imagesampleimplicitlod = 87,
+	spirv_op_imagesampleexplicitlod = 88,
+	spirv_op_negate = 127,
+	spirv_op_fadd = 129,
+	spirv_op_fsub = 131,
+	spirv_op_fmul = 133,
+	spirv_op_fdiv = 136,
+	spirv_op_fmod = 141,
+	spirv_op_vectortimesscalar = 142,
+	spirv_op_vectortimesmatrix = 144,
+	spirv_op_dot = 148,
+	spirv_op_logicalor = 166,
+	spirv_op_logicaland = 167,
+	spirv_op_select = 169,
+	spirv_op_fordequal = 180,
+	spirv_op_fordnotequal = 182,
+	spirv_op_fordlessthan = 184,
+	spirv_op_fordgreaterthan = 186,
+	spirv_op_fordlessthanequal = 188,
+	spirv_op_fordgreaterthanequal = 190,
+	spirv_op_phi = 245,
+	spirv_op_selectionmerge = 247,
+	spirv_op_label = 248,
+	spirv_op_branch = 249,
+	spirv_op_branchconditional = 250,
+	spirv_op_return = 253,
+	spirv_op_returnvalue = 254,
+	spirv_op_typepipestorage = 322,
+	spirv_op_typenamedbarrier = 327,
+	spirv_op_moduleprocessed = 330,
+	spirv_op_executionmodeid = 331,
+	spirv_op_invalid = 0x0000ffff,
+	spirv_op_max = 0x7fffffff
+};
+
+/* may be augmented one day to a ops database */
+u8 *ops_name[0xffff + 1] = {
+	[spirv_op_undef] = "undef",
+	[spirv_op_sourcecontinued] = "source_continued",
+	[spirv_op_source] = "source",
+	[spirv_op_sourceextension] = "source_extension",
+	[spirv_op_name] = "name",
+	[spirv_op_membername] = "member_name",
+	[spirv_op_string] = "string",
+	[spirv_op_line] = "line",
+	[spirv_op_extension] = "extension",
+	[spirv_op_extinstimport] = "ext_inst_import",
+	[spirv_op_extinst] = "ext_inst", 
+	[spirv_op_memorymodel] = "memory_model",
+	[spirv_op_entrypoint] = "entry_point",
+	[spirv_op_executionmode] = "execution_mode",
+	[spirv_op_capability] = "capability",
+	[spirv_op_typevoid] = "type_void",
+	[spirv_op_typebool] = "type_bool",
+	[spirv_op_typeint] = "type_int",
+	[spirv_op_typefloat] = "type_float",
+	[spirv_op_typevector] = "type_vector",
+	[spirv_op_typematrix] = "type_matrix",
+	[spirv_op_typeimage] = "type_image",
+	[spirv_op_typesampler] = "type_sampler",
+	[spirv_op_typesampledimage] = "type_sampled_image",
+	[spirv_op_typearray] = "type_array",
+	[spirv_op_typeruntimearray] = "type_runtime_array",
+	[spirv_op_typestruct] = "type_struct",
+	[spirv_op_typeopaque] = "type_opaque",
+	[spirv_op_typepointer] = "type_pointer",
+	[spirv_op_typefunction] = "type_function",
+	[spirv_op_typeevent] = "type_event",
+	[spirv_op_typedeviceevent] = "type_device_event",
+	[spirv_op_typereserveid] = "type_reserve_id",
+	[spirv_op_typequeue] = "type_queue",
+	[spirv_op_typepipe] = "type_pipe",
+	[spirv_op_typeforwardpointer] = "type_forward_pointer",
+	[spirv_op_constanttrue] = "constant_true",
+	[spirv_op_constantfalse] = "constant_false",
+	[spirv_op_constant] = "constant",
+	[spirv_op_constantcomposite] = "constant_composite",
+	[spirv_op_constantsampler] = "constant_sampler",
+	[spirv_op_constantnull] = "constant_null",
+	[spirv_op_specconstanttrue] = "spec_constant_true",
+	[spirv_op_specconstantfalse] = "spec_constant_false",
+	[spirv_op_specconstant] = "spec_constant",
+	[spirv_op_specconstantcomposite] = "spec_constant_composite",
+	[spirv_op_specconstantop] = "spec_constant_op",
+	[spirv_op_function] = "function",
+	[spirv_op_functionparameter] = "function_parameter",
+	[spirv_op_functionend] = "function_end",
+	[spirv_op_functioncall] = "function_call",
+	[spirv_op_variable] = "variable",
+	[spirv_op_load] = "load",
+	[spirv_op_store] = "store",
+	[spirv_op_accesschain] = "accesschain",
+	[spirv_op_decorate] = "decorate",
+	[spirv_op_memberdecorate] = "member_decorate",
+	[spirv_op_decorationgroup] = "decoration_group",
+	[spirv_op_groupdecorate] = "group_decorate",
+	[spirv_op_groupmemberdecorate] = "group_member_decorate",
+	[spirv_op_vectorshuffle] = "vector_shuffle",
+	[spirv_op_compositeconstruct] = "composite_construct",
+	[spirv_op_compositeextract] = "composite_extract",
+	[spirv_op_compositeinsert] = "composite_insert",
+	[spirv_op_sampledimage] = "sampled_image",
+	[spirv_op_imagesampleimplicitlod] = "image_sample_implicit_lod",
+	[spirv_op_imagesampleexplicitlod] = "image_sample_explicit_lod",
+	[spirv_op_negate] = "negate",
+	[spirv_op_fadd] = "fadd",
+	[spirv_op_fsub] = "fsub",
+	[spirv_op_fmul] = "fmul",
+	[spirv_op_fdiv] = "fdiv",
+	[spirv_op_fmod] = "fmod",
+	[spirv_op_vectortimesscalar] = "vector_times_scalar",
+	[spirv_op_vectortimesmatrix] = "vector_times_matrix",
+	[spirv_op_dot] = "dot",
+	[spirv_op_logicalor] = "logical_or",
+	[spirv_op_logicaland] = "logical_and",
+	[spirv_op_select] = "select", 
+	[spirv_op_fordequal] = "ford_equal",
+	[spirv_op_fordnotequal] = "ford_not_equal",
+	[spirv_op_fordlessthan] = "ford_less_than",
+	[spirv_op_fordgreaterthan] = "ford_greater_than",
+	[spirv_op_fordlessthanequal] = "ford_less_than_equal",
+	[spirv_op_fordgreaterthanequal] = "ford_greater_than_equal",
+	[spirv_op_phi] = "phi",
+	[spirv_op_selectionmerge] = "selection_merge",
+	[spirv_op_label] = "label",
+	[spirv_op_branch] = "branch",
+	[spirv_op_branchconditional] = "branch_conditional",
+	[spirv_op_return] = "return",
+	[spirv_op_returnvalue] = "return_value",
+	[spirv_op_typepipestorage] = "type_pipe_storage",
+	[spirv_op_typenamedbarrier] = "type_named_barrier",
+	[spirv_op_moduleprocessed] = "module_processed",
+	[spirv_op_executionmodeid] = "execution_mode_id",
+	[spirv_op_invalid] = "invalid"
+};
+
+struct hdr {
+	u32 magic_num;
+	u32 version;
+	u32 generator;
+	u32 bound;
+	u32 reserved;
+};
+
+static u8 depth;
+
+static u32 op;
+#define op_ws_n (op >> 16)
+#define op_ws_last (op_ws_n - 1) /* last op w idx */
+#define op_ws_end op_ws_n /* idx right after the last op w */
+#define op_num (op & 0x0000ffff)
+
+static u32 opds[0x0000ffff + 1];
+#define opds_n (op_ws_n - 1)
+#define opds_last (opds_n - 1) /* last opd w idx */
+#define opds_end opds_n /* w idx right after the last opd w*/
+
+#define op_name ops_name[op_num]
+
+#define OK 0
+#define END 1
+#define ERR 2
+static u8 read_whole_op(void)
+{
+	size_t read_bytes_n;
+	read_bytes_n = fread(&op, sizeof(op), 1, stdin);
+	if (read_bytes_n != sizeof(op))	{
+		if (feof(stdin) != 0)
+			return END;
+		if (ferror(stdin) != 0)
+			return ERR;
+	}
+
+	read_bytes_n = fread(opds, sizeof(u32), opds_n, stdin);
+	if (read_bytes_n != (sizeof(u32) * opds_n)) {	
+		if (feof(stdin) != 0)
+			return END;
+		if (ferror(stdin) != 0)
+			return ERR;
+	}
+}
+
+static void hdr_out(void)
+{
+	struct hdr hdr;
+
+	(void)fread(&hdr, sizeof(hdr), 1, stdin);
+	if (feof(stdin) != 0 || ferror(stdin) != 0)
+		exit(0); 
+
+	out("// magic number = 0x%08x (%s)\n", hdr.magic_num, hdr.magic_num == 0x07230203 ? "good" : "bad");
+	out("// version = 0x%08x\n", hdr.version);
+	out("// generator = 0x%08x\n", hdr.generator);
+	out("// bound = %u\n", hdr.bound);
+	out("// reserved = 0x%08x\n", hdr.reserved);
+}
+
+static void layout_caps_out(void)
+{
+	u16 section_breaking_op_num;
+
+	section_breaking_op_num = spirv_op_invalid;
+	out("\n// section start: capabilities\n");
+	loop {
+		u8 r;
+
+		if (op_num != spirv_op_capability) {
+			section_breaking_op_num = op_num;
+			break;
+		}
+
+		out("%s capability=0x%08x\n", op_name, opds[0]);
+
+		r = read_whole_op();
+		if (r != OK)
+			break;
+	}
+	if (section_breaking_op_num == spirv_op_invalid) {
+		out("// section end: capabilities\n");
+		exit(0);
+	}
+	out("// section end: capabilities, breaking opcode=%s(%u)\n", ops_name[section_breaking_op_num], section_breaking_op_num);
+}
+
+/* return the sz of the litstr as a n of ws */ 
+static u16 litstr_out(u32 *p)
+{
+	u8 b[4];
+	u32 *w;
+	u16 ws_n;
+
+	w = (u32*)b;
+	out("\"");
+	ws_n = 1;
+	loop {
+		u8 i;
+
+		*w = *p;
+		i = 0;
+		loop {
+			if (i == 4)
+				break;
+
+			if (b[i] == 0) { /* 0 terminating char */
+				out("\"");
+				return ws_n;
+			}
+
+			(void)fwrite(&b[i], 1, 1, stdout);
+
+			++i;
+		}
+		++p;
+		++ws_n;
+	}
+}
+
+static void layout_extns_out(void)
+{
+	u16 section_breaking_op_num;
+
+	section_breaking_op_num = spirv_op_invalid;
+	out("\n// section start: extensions\n");
+	loop {
+		u8 r;
+
+		if (op_num != spirv_op_extension) {
+			section_breaking_op_num = op_num;
+			break;
+		}
+
+		out("%s ", op_name);
+		(void)litstr_out(opds);
+		out("\n");
+
+		r = read_whole_op();
+		if (r != OK)
+			break;
+	}
+	if (section_breaking_op_num == spirv_op_invalid) {
+		out("// section end: extensions\n");
+		exit(0);
+	}
+	out("// section end: extensions, breaking opcoder=%s(%u)\n", ops_name[section_breaking_op_num], section_breaking_op_num);
+}
+
+static void layout_extinstimports_out(void)
+{
+	u16 section_breaking_op_num;
+
+	section_breaking_op_num = spirv_op_invalid;
+	out("\n// section start: extended set of instructions imports\n");
+	loop {
+		u8 r;
+
+		if (op_num != spirv_op_extinstimport) {
+			section_breaking_op_num = op_num;
+			break;
+		}
+
+		out("%%%u = %s name=", opds[0], op_name);
+		litstr_out(opds + 1);
+		out("\n");
+
+		r = read_whole_op();
+		if (r != OK)
+			break;
+	}
+	if (section_breaking_op_num == spirv_op_invalid) {
+		out("// section end: extended set of instructions imports\n");
+		exit(0);
+	}
+	out("// section end: extended set of instructions imports, breaking opcode=%s(%u)\n", ops_name[section_breaking_op_num], section_breaking_op_num);
+}
+
+static u8 *addressing_model_to_str(u32 w)
+{
+	switch (w) {
+	case 0:
+		return "logical";
+	case 1:
+		return "physical32";
+	case 2:
+		return "physical64";
+	default:
+		return "unknown_addressing_model_code";
+	}
+}
+
+static u8 *memory_model_to_str(u32 w)
+{
+	switch (w) {
+	case 0:
+		return "simple";
+	case 1:
+		return "glsl450";
+	case 2:
+		return "opencl";
+	case 3:
+		return "vulkan";
+	default:
+		return "unkwnown_memory_model_code";
+	}
+}
+
+static void layout_memorymodel_out(void)
+{
+	u8 r;
+
+	out("\n// the only memory model instruction, if one, should be here\n");
+	if (op_num == spirv_op_memorymodel)
+		out("%s addressing_model=%s memory_model=%s\n", op_name, addressing_model_to_str(opds[0]), memory_model_to_str(opds[1]));
+
+	r = read_whole_op();
+	if (r != OK)
+		exit(0);
+}
+
+static u8 *execution_model_to_str(u32 w)
+{
+	switch (w) {
+	case 0:
+		return "vertex";
+	case 1:
+		return "tessellation_control";
+	case 2:
+		return "tessellation_evaluation";
+	case 3:
+		return "geometry";
+	case 4:
+		return "fragment";
+	case 5:
+		return "gl_compute";
+	case 6:
+		return "kernel";
+	default:
+		return "unkwown_execution_model_code";
+	}
+}
+
+static void layout_entrypoints_out(void)
+{
+	u16 section_breaking_op_num;
+
+	section_breaking_op_num = spirv_op_invalid;
+	out("\n// section start: entry points\n");
+	loop {
+		u8 r;
+		u16 name_ws_n;
+		u16 intf_ws_n;
+
+		if (op_num != spirv_op_entrypoint) {
+			section_breaking_op_num = op_num;
+			break;
+		}
+
+		out("%s execution_model=%s entry_point=%%%u name=", op_name, execution_model_to_str(opds[0]), opds[1]);
+		name_ws_n = litstr_out(opds + 2);
+
+		intf_ws_n = op_ws_n - (3 + name_ws_n);
+		if (intf_ws_n != 0) {
+			u16 i;
+
+			i = 0;
+			loop {
+				if (i == intf_ws_n)
+					break;
+				out(" interfaces[%u]=%%%u", i, opds[2 + name_ws_n + i]);
+				++i;
+			}
+		}
+		out("\n");
+
+		r = read_whole_op();
+		if (r != OK)
+			break;
+	}
+	if (section_breaking_op_num == spirv_op_invalid) {
+		out("// section end: entry points\n");
+		exit(0);
+	}
+	out("// section end: entry points, breaking opcode=%s(%u)\n", ops_name[section_breaking_op_num], section_breaking_op_num);
+}
+
+static void op_executionmode_out(void)
+{
+	u16 execution_mode_lits_ws_n;
+	u16 i;
+
+	/*
+	 * we don't decode the execution modes: we'll add the ones we are
+	 * interested in
+	 */
+	out("%s entry_point=%%%u mode=0x%08x\n", op_name, opds[0], opds[1]);
+
+	execution_mode_lits_ws_n = op_ws_n - 3;
+	if (execution_mode_lits_ws_n == 0 )
+		return;
+
+	i = 0;
+	loop {
+		if (i == execution_mode_lits_ws_n)
+			break;
+		out(" 0x%08x", opds[2 + i]);
+		++i;
+	}
+}
+
+static void op_executionmodeid_out(void)
+{	
+	u16 execution_mode_ids_ws_n;
+	u16 i;
+
+	out("%s %%%u 0x%08x ", op_name, opds[0], opds[1]);
+
+	execution_mode_ids_ws_n = op_ws_n - 3;
+	if (execution_mode_ids_ws_n == 0)
+		return;
+
+	i = 0;
+	loop {
+		if (i == execution_mode_ids_ws_n)
+			break;
+		out(" %%%u", opds[2 + i]);
+		++i;
+	}
+}
+
+static void layout_executionmodes_out(void)
+{
+	u16 section_breaking_op_num;
+
+	section_breaking_op_num = spirv_op_invalid;
+	out("\n// section start: execution modes\n");
+	loop {
+		u8 r;
+
+		if (op_num == spirv_op_executionmode) {
+			op_executionmode_out();
+		} else if (op_num == spirv_op_executionmodeid) {
+			op_executionmodeid_out();
+		} else {
+			section_breaking_op_num = op_num;
+			break;
+		}
+
+		r = read_whole_op();
+		if (r != OK)
+			break;
+	}
+	if (section_breaking_op_num == spirv_op_invalid) {
+		out("// section end: execution modes\n");
+		exit(0);
+	}
+	out("// section end: execution modes, breaking opcoder=%s(%u)\n", ops_name[section_breaking_op_num], section_breaking_op_num);
+}
+
+static u8 *src_lang_str(u32 w)
+{
+	switch (w) {
+	case 0:
+		return "unknown";
+	case 1:
+		return "essl";
+	case 2:
+		return "glsl";
+	case 3:
+		return "opencl_c";
+	case 4:
+		return "opencl_cpp";
+	case 5:
+		return "hlsl";
+	default:
+		return "implicit_unknow";
+	}
+}
+
+static void op_sourcecontinued_out(void)
+{
+	out("%s ", op_name);
+	(void)litstr_out(opds);
+	out("\n");
+}
+
+static void op_source_out(void)
+{
+	out("5s %s 0x%08x", op_name, src_lang_str(opds[0]), opds[1]);
+	if (op_ws_n > 3) {
+		out(" %%%u", opds[2]);
+
+		if (op_ws_n > 4) {
+			out(" ");
+			(void)litstr_out(opds + 3);
+		}
+	}
+	out("\n");
+}
+
+static void op_sourceextension_out(void)
+{
+	out("%s ", op_name);
+	(void)litstr_out(opds);
+	out("\n");
+}
+
+static void op_string_out(void)
+{
+	out("%s %%%u", op_name, opds[0]);
+	(void)litstr_out(opds + 1);
+	out("\n");
+}
+
+static bool layout_debug_section_0(void)
+{
+	u16 section_breaking_op_num;
+
+	section_breaking_op_num = spirv_op_invalid;
+	out("\n// debug first subsection start\n");
+	loop {
+		u8 r;
+
+		if (op_num == spirv_op_sourcecontinued) {
+			op_sourcecontinued_out();
+		} else if (op_num == spirv_op_source) {
+			op_source_out();
+		} else if (op_num == spirv_op_sourceextension) {
+			op_sourceextension_out();
+		} else if (op_num == spirv_op_string) {
+			op_string_out();
+		} else {
+			section_breaking_op_num = op_num;
+			break;
+		}
+
+		r = read_whole_op();
+		if (r != OK)
+			break;
+	}
+	if (section_breaking_op_num == spirv_op_invalid) {
+		out("// debug first subsection end\n");
+		return true;
+	}
+	out("// debug first subsection end, breaking opcode=%s(%u)\n", ops_name[section_breaking_op_num], section_breaking_op_num);
+	return false;
+}
+
+static void op_name_out(void)
+{
+	out("%s target=%%%u name=", op_name, opds[0]);
+	(void)litstr_out(opds + 1);
+	out("\n");
+}
+
+static void op_membername_out(void)
+{
+	out("%s type=%%%u member=%u name=", op_name, opds[0], opds[1]);
+	(void)litstr_out(opds + 2);
+	out("\n");
+}
+
+static bool layout_debug_section_1(void)
+{
+	u16 section_breaking_op_num;
+
+	section_breaking_op_num = spirv_op_invalid;
+	out("\n// debug second subsection start\n");
+	loop {
+		u8 r;
+
+		if (op_num == spirv_op_name) {
+			op_name_out();
+		} else if (op_num == spirv_op_membername) {
+			op_membername_out();
+		} else {
+			section_breaking_op_num = op_num;
+			break;
+		}
+
+		r = read_whole_op();
+		if (r != OK)
+			break;
+	}
+	if (section_breaking_op_num == spirv_op_invalid) {
+		out("// debug second subsection end\n");
+		return true;
+	}
+	out("// debug second subsection end, breaking opcode=%s(%u)\n", ops_name[section_breaking_op_num], section_breaking_op_num);
+	return false;
+}
+
+static bool layout_debug_section_2(void)
+{
+	u16 section_breaking_op_num;
+
+	section_breaking_op_num = spirv_op_invalid;
+	out("\n// debug third subsection start\n");
+	loop {
+		u8 r;
+
+		if (op_num != spirv_op_moduleprocessed) {
+			section_breaking_op_num = op_num;
+			break;
+		}
+		out("%s ", op_name);
+		(void)litstr_out(opds);
+		out("\n");
+
+		r = read_whole_op();
+		if (r != OK)
+			break;
+	}
+	if (section_breaking_op_num == spirv_op_invalid) {
+		out("// debug third subsection end\n");
+		return true;
+	}
+	out("// debug third subsection end, breaking opcode=%s(%u)\n", ops_name[section_breaking_op_num], section_breaking_op_num);
+	return false;
+}
+
+static void layout_debug_out(void)
+{
+	bool do_exit;
+
+	out("\n// section start: debug\n");
+
+	do_exit = layout_debug_section_0();
+	if (do_exit)
+		goto exit;
+
+	do_exit = layout_debug_section_1();
+	if (do_exit)
+		goto exit;
+
+	do_exit = layout_debug_section_2();
+
+exit:
+	out("// section end: debug\n");
+	if (do_exit)
+		exit(0);
+}
+
+static void decoration_builtin_out(u16 i)
+{
+	u8 *str;
+
+	switch (opds[i + 1]) {
+	case 0:
+		str = "position";
+		break;
+	case 1:
+		str = "point_size";
+		break;	
+	case 3:
+		str = "clip_distance";
+		break;	
+	case 4:
+		str = "cull_distance";
+		break;
+	case 5:
+		str = "vertex_id";
+		break;
+	case 6:
+		str = "instance_id";
+		break;
+	case 7:
+		str = "primitive_id";
+		break;
+	case 8:
+		str = "invocation_id";
+		break;
+	case 9:
+		str = "layer";
+		break;
+	case 10:
+		str = "viewport_idx";
+		break;
+	case 11:
+		str = "tess_level_outer";
+		break;
+	case 12:
+		str = "tess_level_inner";
+		break;
+	case 13:
+		str = "tess_coord";
+		break;
+	case 14:
+		str = "patch_vertices";
+		break;
+	case 15:
+		str = "frag_coord";
+		break;
+	case 16:
+		str = "point_coord";
+		break;
+	case 17:
+		str = "front_facing";
+		break;
+	case 18:
+		str = "sample_id";
+		break;
+	case 19:
+		str = "sample_position";
+		break;
+	case 20:
+		str = "sample_mask";
+		break;
+	case 22:
+		str = "frag_depth";
+		break;
+	case 23:
+		str = "helper_invocation";
+		break;
+	case 24:
+		str = "num_workgroups";
+		break;
+	case 25:
+		str = "workgroup_size";
+		break;
+	case 26:
+		str = "workgroup_id";
+		break;
+	case 27:
+		str = "local_invocation_id";
+		break;
+	case 28:
+		str = "global_invocation_id";
+		break;
+	case 29:
+		str = "global_invocation_idx";
+		break;
+	case 30:
+		str = "work_dim";
+		break;
+	case 31:
+		str = "global_size";
+		break;
+	case 32:
+		str = "enqueue_workgroup_size";
+		break;
+	case 33:
+		str = "global_offset";
+		break;
+	case 34:
+		str = "global_linear_id";
+		break;
+	case 36:
+		str = "subgroup_size";
+		break;
+	case 37:
+		str = "subgroup_max_size";
+		break;
+	case 38:
+		str = "num_subgroups";
+		break;
+	case 39:
+		str = "num_enqueued_subgroups";
+		break;
+	case 40:
+		str = "subgroup_id";
+		break;
+	case 41:
+		str = "subgroup_local_invocation_id";
+		break;
+	case 42:
+		str = "vertex_index";
+		break;
+	case 43:
+		str = "instance_index";
+		break;
+	case 4424:
+		str = "base_vertex";
+		break;
+	case 4425:
+		str = "base_instance";
+		break;
+	case 4426:
+		str = "draw_index";
+		break;
+	case 4438:
+		str = "device_index";
+		break;
+	case 4440:
+		str = "view_index";
+		break;
+	default:
+		str = "unknown_builtin_code";
+		break;
+	}
+
+	out(" builtin %s", str);
+}
+
+static u8 *scopeid_to_str(u32 w)
+{
+	switch (w) {
+	case 0:
+		return "cross_device";
+	case 1:
+		return "device";
+	case 2:
+		return "workgroup";
+	case 3:
+		return "subgroup";
+	case 4:
+		return "invocation";
+	case 5:
+		return "queue_family";
+	default:
+		return "unknow_scope_id_code";
+	}
+}
+
+static u8 *fproundingmode_to_str(u32 w)
+{
+	switch (w) {
+	case 0:
+		return "rte";
+	case 1:
+		return "rtz";
+	case 2:
+		return "rtp";
+	case 3:
+		return "rtn";
+	default:
+		return "unknown_fp_rounding_mode_code";
+	}
+}
+
+#define flag(val, str) \
+if ((w & val) != 0) { \
+	if (!first) \
+		strcat(b, "|"); \
+	strcat(b, str); \
+	first = false; \
+}
+static u8 *fpfastmathmode_to_str(u32 w)
+{
+	static u8 b[256];
+	bool first;
+
+	if (w == 0)
+		return "none";
+
+	b[0] = 0;
+	first = true;
+
+	flag(0x00000001, "not_nan")
+	flag(0x00000002, "not_inf")
+	flag(0x00000004, "nsz")
+	flag(0x00000008, "allow_recip")
+	flag(0x00000010, "fast")
+	flag(0xffffffe0, "unknown_fp_fast_math_mode_code")
+}
+#undef flag
+
+static u8 *funcparamattr_to_str(u32 w)
+{
+	switch (w) {
+	case 0:
+		return "z_ext";
+	case 1:
+		return "s_ext";
+	case 2:
+		return "by_val";
+	case 3:
+		return "s_ret";
+	case 4:
+		return "no_alias";
+	case 5:
+		return "no_capture";
+	case 6:
+		return "no_write";
+	case 7:
+		return "no_read_write";
+	defaut:
+		return "unknown_func_param_attr_code";
+	}
+}
+
+
+/* i is the start idx in the opds array of the decoration specifications */
+static void decoration_out(u16 i)
+{
+	switch (opds[i]) {
+	case 0:
+		out(" relaxedr_|precision");
+		break;
+	case 1:
+		out(" spec_id");
+		break;
+	case 2:
+		out(" block");
+		break;
+	case 3:
+		out(" buffer_block");
+		break;
+	case 4:
+		out(" row_major");
+		break;
+	case 5:
+		out(" col_major");
+		break;
+	case 6:
+		out(" array_stride %u", opds[i + 1]);
+		break;
+	case 7:
+		out(" matrix_stride %u", opds[i + 1]);
+		break;
+	case 8:
+		out(" glsl_shared");
+		break;
+	case 9:
+		out(" glsl_packed");
+		break;
+	case 10:
+		out(" c_packed");
+		break;
+	case 11:
+		decoration_builtin_out(i);
+		break;
+	case 13:
+		out(" no_perspective");
+		break;
+	case 14:
+		out(" flat");
+		break;
+	case 15:
+		out(" patch");
+		break;
+	case 16:
+		out(" centroid");
+		break;
+	case 17:
+		out(" sample");
+		break;
+	case 18:
+		out(" invariant");
+		break;
+	case 19:
+		out(" restrict");
+		break;
+	case 20:
+		out(" aliased");
+		break;
+	case 21:
+		out(" volatile");
+		break;
+	case 22:
+		out(" constant");
+		break;
+	case 23:
+		out(" coherent");
+		break;
+	case 24:
+		out(" nonwritable");
+		break;
+	case 25:
+		out(" nonreadable");
+		break;
+	case 26:
+		out(" uniform");
+		break;
+	case 27:
+		out(" uniform_id %s", scopeid_to_str(opds[i + 1]));
+		break;
+	case 28:
+		out(" saturated_conversion");
+		break;
+	case 29:
+		out(" stream %u", opds[i + 1]);
+		break;
+	case 30:
+		out(" location %u", opds[i + 1]);
+		break;
+	case 31:
+		out(" component %u", opds[i + 1]);
+		break;
+	case 32:
+		out(" index %u", opds[i + 1]);
+		break;
+	case 33:
+		out(" binding %u", opds[i + 1]);
+		break;
+	case 34:
+		out(" descriptor_set %u", opds[i + 1]);
+		break;
+	case 35:
+		out(" offset %u", opds[i + 1]);
+		break;
+	case 36:
+		out(" xfb_buffer %u", opds[i + 1]);
+		break;
+	case 37:
+		out(" xfb_stride %u", opds[i + 1]);
+		break;
+	case 38:
+		out(" func_param_attr %s", funcparamattr_to_str(opds[i + 1]));
+		break;
+	case 39:
+		out(" fp_rounding_mode %s", fproundingmode_to_str(opds[i + 1]));
+		break;
+	case 40:
+		out(" fp_fast_math_mode %s", fpfastmathmode_to_str(opds[i + 1]));
+		break;
+	case 41:
+		out(" linkage_attributes (...)");
+		break;
+	case 42:
+		out(" no_contraction");
+		break;
+	case 43:
+		out(" input_attachment_index %u", opds[i + 1]);
+		break;
+	case 44:
+		out(" alignment %u", opds[i + 1]);
+		break;
+
+	case 45:
+		out(" max_byte_offset %u", opds[i + 1]);
+		break;
+	case 46:
+		out(" alignment_id %%%u", opds[i + 1]);
+		break;
+	case 47:
+		out(" max_byte_offset_id %%%u", opds[i + 1]);
+		break;
+	default:
+		out(" decoration_not_handled(%u operands)", op_ws_n - i - 1);
+		break;
+	}
+}
+
+static void op_decorate_out(void)
+{
+	out("%s %%%u", op_name, opds[0]);
+	decoration_out(1);
+	out("\n");
+}
+
+static void op_memberdecorate_out(void)
+{
+	out("%s %%%u %u", op_name, opds[0], opds[1]);
+	decoration_out(2);
+	out("\n");
+}
+
+static void layout_annotation_out(void)
+{
+	u16 section_breaking_op_num;
+
+	section_breaking_op_num = spirv_op_invalid;
+	out("\n// section start: annotations\n");
+	loop {
+		u8 r;
+
+		if (op_num == spirv_op_decorate) {
+			op_decorate_out();
+		} else if (op_num == spirv_op_memberdecorate) {
+			op_memberdecorate_out();
+		} else if (op_num == spirv_op_groupdecorate)  {
+			out("%s(deprecated, SKIPPING)\n", op_name);
+		} else if (op_num == spirv_op_groupmemberdecorate) {
+			out("%s(deprecated, SKIPPING)\n", op_name);
+		} else if (op_num == spirv_op_decorationgroup) {
+			out("%s(deprecated, SKIPPING)\n", op_name);
+		} else {
+			section_breaking_op_num = op_num;
+			break;
+		}
+
+		r = read_whole_op();
+		if (r != OK)
+			break;
+	}
+	if (section_breaking_op_num == spirv_op_invalid) {
+		out("// section end: annotations\n");
+		exit(0);
+	}
+	out("// section end: annotations, breaking opcode=%s(%u)\n", ops_name[section_breaking_op_num], section_breaking_op_num);
+}
+
+static u8 *signedness_to_str(u32 w)
+{
+	switch (w) {
+	case 0:
+		return "unsigned";
+	case 1:
+		return "signed";
+	default:
+		return "unknown_signedness_code";
+	}
+}
+
+static u8 *dim_to_str(u32 w)
+{
+	switch (w) {
+	case 0:
+		return "1d";
+	case 1:
+		return "2d";
+	case 2:
+		return "3d";
+	case 3:
+		return "cube";
+	case 4:
+		return "rect";
+	case 5:
+		return "buffer";
+	case 6:
+		return "subpass_data";
+	default:
+		return "unknown_dimension_code";
+	}
+}
+
+static u8 *depth_to_str(u32 w)
+{
+	switch (w) {
+	case 0:
+		return "no_depth_image";
+	case 1:
+		return "depth_image";
+	case 2:
+		return "no_depth_information";
+	default:
+		return "unknown_depth_code";
+	}
+}
+
+static u8 *arrayed_to_str(u32 w)
+{
+	switch (w) {
+	case 0:
+		return "non_arrayed_content";
+	case 1:
+		return "arrayed_content";
+	default:
+		return "unknown_arrayed_code";
+	}
+}
+
+static u8 *multisample_to_str(u32 w)
+{
+	switch (w) {
+	case 0:
+		return "single_sampled";
+	case 1:
+		return "multi_sampled";
+	default:
+		return "unknown_multisample_code";
+	}
+}
+
+static u8 *sampled_to_str(u32 w)
+{
+	switch (w) {
+	case 0:
+		return "runtime_known";
+	case 1:
+		return "sampler";
+	case 2:
+		return "no_sampler";
+	default:
+		return "unknown_sampled_code";
+	}
+}
+
+static u8 *img_fmt_to_str(u32 w)
+{
+	switch (w) {
+	case 0:
+		return "unkown";
+	case 1:
+		return "rgba32f";
+	case 2:
+		return "rgba16f";
+	case 3:
+		return "r32f";
+	case 4:
+		return "rgba8";
+	case 5:
+		return "rgba8snorm";
+	case 6:
+		return "rg32f";
+	case 7:
+		return "rg16f";
+	case 8:
+		return "r11g11b10f";
+	case 9:
+		return "r16f";
+	case 10:
+		return "rgba16";
+	case 11:
+		return "rgb10a2";
+	case 12:
+		return "rg16";
+	case 13:
+		return "rg8";
+	case 14:
+		return "r16";
+	case 15:
+		return "r8";
+	case 16:
+		return "rgba16snorm";
+	case 17:
+		return "rg16snorm";
+	case 18:
+		return "rg8snorm";
+	case 19:
+		return "r16snorm";
+	case 20:
+		return "r8snorm";
+	case 21:
+		return "rgba32i";
+	case 22:
+		return "rgba16i";
+	case 23:
+		return "rgba8i";
+	case 24:
+		return "r32i";
+	case 25:
+		return "rg32i";
+	case 26:
+		return "rg16i";
+	case 27:
+		return "rg8i";
+	case 28:
+		return "r16i";
+	case 29:
+		return "r8i";
+	case 30:
+		return "rgba32ui";
+	case 31:
+		return "rgba16ui";
+	case 32:
+		return "rgba8ui";
+	case 33:
+		return "r32ui";
+	case 34:
+		return "rgb10a2ui";
+	case 35:
+		return "rg32ui";
+	case 36:
+		return "rg16ui";
+	case 37:
+		return "rg8ui";
+	case 38:
+		return "r16ui";
+	case 39:
+		return "r8ui";
+	default:
+		return "unknown_image_format_code";
+	}
+}
+
+static u8 *access_qualifier_to_str(u32 w)
+{
+	switch (w) {
+	case 0:
+		return "read_only";
+	case 1:
+		return "write_only";
+	case 2:
+		return "read_write";
+	default:
+		return "unknown_access_qualifier_code";
+	}
+}
+
+static void op_typeimage_out(void)
+{
+	out("%%%u = %s sampled_type=%%%u dim=%s depth=%s arrayed=%s multisample=%s sampled=%s image_format=%s", opds[0], op_name, opds[1], dim_to_str(opds[2]), depth_to_str(opds[3]), arrayed_to_str(opds[4]), multisample_to_str(opds[5]), sampled_to_str(opds[6]), img_fmt_to_str(opds[7]));
+	if (op_ws_n > 9) /* have access qualifier */
+		out(" access_qualifier=%s\n", access_qualifier_to_str(opds[8]));
+	else
+		out("\n");
+}
+
+static void op_typestruct_out(void)
+{
+	u16 i;
+
+	out("%%%u = %s", opds[0], op_name);
+	i = 1;
+	loop  {
+		if (i > opds_last)
+			break;
+		out(" member_%u=%%%u", i - 1, opds[i]);
+		++i;
+	}
+	out("\n");
+}
+
+static u8 *storage_class_to_str(u32 w)
+{
+	switch (w) {
+	case 0:
+		return "uniform_constant";
+	case 1:
+		return "input";
+	case 2:
+		return "uniform";
+	case 3:
+		return "output";
+	case 4:
+		return "workgroup";
+	case 5:
+		return "cross_workgroup";
+	case 6:
+		return "private";
+	case 7:
+		return "function";
+	case 8:
+		return "generic";
+	case 9:
+		return "push_constant";
+	case 10:
+		return "atomic_counter";
+	case 11:
+		return "image";
+	case 12:
+		return "storage_buffer";
+	case 5349:
+		return "physical_storage_buffer";
+	default:
+		return "unknown_storage_class_code";
+	}
+}
+
+static void op_typefunction_out(void)
+{
+	u16 i;
+
+	out("%%%u = %s return_type=%%%u", opds[0], op_name, opds[1]);
+
+	i = 2;
+	loop {
+		if (i > opds_last)
+			break;
+		out(" parameters[%u]=%%%u", i - 2, opds[i]);
+		++i;
+	}	
+	out("\n");
+}
+
+static void op_constant_out(void)
+{
+	u16 i;
+
+	out("%%%u = %s type=%%%u", opds[1], op_name, opds[0]);
+
+	i = 2;
+	loop {
+		if (i > opds_last)
+			break;
+		out(" values[%u]=0x%08x", i - 2, opds[i]);
+		++i;
+	}
+	out("\n");
+}
+
+static void op_specconstant_out(void)
+{
+	u16 i;
+
+	out("%%%u = %s type=%%%u", opds[1], op_name, opds[0]);
+
+	i = 2;
+	loop {
+		if (i > opds_last)
+			break;
+		out(" values[%u]=0x%08x", i - 2, opds[i]);
+		++i;
+	}
+	out("\n");
+}
+
+static void op_constantcomposite_out(void)
+{
+	u16 i;
+
+	out("%%%u = %s type=%%%u", opds[1], op_name, opds[0]);
+
+	i = 2;
+	loop {
+		if (i > opds_last)
+			break;
+		out(" constituents[%u]=%%%u", i - 2, opds[i]);
+		++i;
+	}
+	out("\n");
+}
+
+static void op_specconstantcomposite_out(void)
+{
+	u16 i;
+
+	out("%%%u = %s type=%%%u", opds[1], op_name, opds[0]);
+
+	i = 2;
+	loop {
+		if (i > opds_last)
+			break;
+		out(" constituents[%u]=%%%u", i - 2, opds[i]);
+		++i;
+	}
+	out("\n");
+}
+
+static u8 *sampler_addressing_mode_to_str(u32 w)
+{
+	switch (w) {
+	case 0:
+		return "none";
+	case 1:
+		return "clamp_to_edge";
+	case 2:
+		return "clamp";
+	case 3:
+		return "repeat";
+	case 4:
+		return "repeat_mirrored";
+	default:
+		return "unknown_sampler_addressing_mode_code";
+	}
+}
+
+static u8 *sampler_filter_mode_to_str(u32 w)
+{
+	switch (w) {
+	case 0:
+		return "nearest";
+	case 1:
+		return "linear";
+	default:
+		return "unknown_sampler_filter_mode_code";
+	}
+}
+
+static void op_constantsampler_out(void)
+{
+	u8 *param;
+
+	switch (opds[3]) {
+	case 0:
+		param = "non_normalized";
+		break;
+	case 1:
+		param = "normalized";
+		break;
+	default:
+		param = "unknown_param_code";
+		break;
+	}
+	out("%%%u = 5s type=%%%u sampler_addressing_mode=%s param=%s sampler_filter_mode=%s", opds[1], op_name, opds[0], sampler_addressing_mode_to_str(opds[2]), param, sampler_filter_mode_to_str(opds[3]));
+}
+
+static void op_specconstantop_out(void)
+{
+	u16 i;
+
+	out("%%%u = %s type=%%%u opcode=%u", opds[1], op_name, opds[0], opds[2]);
+
+	i = 3;
+	loop {
+		if (i > opds_last)
+			break;
+		out(" operands[%u]=0x%08x", i - 3, opds[i]);
+		++i;
+	}
+	out("\n");
+}
+
+static void op_variable_out(bool depth)
+{
+	if (depth) {
+		out_depth("%%%u = %s pointer_type=%%%u storage_class=%s", opds[1], op_name, opds[0], storage_class_to_str(opds[2]));
+	} else {
+		out("%%%u = %s pointer_type=%%%u storage_class=%s", opds[1], op_name, opds[0], storage_class_to_str(opds[2]));
+	}
+	if (op_ws_n > 4)
+		out(" initializer=%%%u", opds[3]);
+	out("\n");
+}
+
+static void layout_nonfuncdecls_out(void)
+{
+	u16 section_breaking_op_num;
+
+	section_breaking_op_num = spirv_op_invalid;
+	out("\n// section start: non function declarations\n");
+	loop {
+		u8 r;
+
+		switch (op_num) {
+		case spirv_op_line:
+			out("%s %%%u line=%u column=%u\n", op_name, opds[0], opds[1], opds[2]);
+			break;
+		/* types start -----------------------------------------------*/
+		case spirv_op_typevoid:
+			out("%%%u = %s\n", opds[0], op_name);
+			break;
+		case spirv_op_typebool:
+			out("%%%u = %s\n", opds[0], op_name);
+			break;
+		case spirv_op_typeint:
+			out("%%%u = %s width=%u signedness=%s\n", opds[0], op_name, opds[1], signedness_to_str(opds[2]));
+			break;
+		case spirv_op_typefloat:
+			out("%%%u = %s width=%u\n", opds[0], op_name, opds[1]);
+			break;
+		case spirv_op_typevector:
+			out("%%%u = %s component_type=%%%u component_count=%u\n", opds[0], op_name, opds[1], opds[2]);
+			break;
+		case spirv_op_typematrix:
+			out("%%%u = %s column_type=%%%u column_count=%u\n", opds[0], op_name, opds[1], opds[2]);
+			break;
+		case spirv_op_typeimage:
+			op_typeimage_out();
+			break;
+		case spirv_op_typesampler:
+			out("%%%u = %s\n", opds[0], op_name);
+			break;
+		case spirv_op_typesampledimage:
+			out("%%%u = %s image_type=%%%u\n", opds[0], op_name, opds[1]);
+			break;
+		case spirv_op_typearray:
+			out("%%%u = %s element_type=%%%u length=%%%u\n", opds[0], op_name, opds[1]);
+			break;
+		case spirv_op_typeruntimearray:
+			out("%%%u = %s element_type=%%%u\n", opds[0], op_name, opds[1]);
+			break;
+		case spirv_op_typestruct:
+			op_typestruct_out();
+			break;
+		case spirv_op_typeopaque:
+			out("%%%u = %s ", opds[0], op_name);
+			(void)litstr_out(opds + 1);
+			out("\n");
+			break;
+		case spirv_op_typepointer:
+			out("%%%u = %s storage_class=%s type=%%%u\n", opds[0], op_name, storage_class_to_str(opds[1]), opds[2]);
+			break;
+		case spirv_op_typefunction:
+			op_typefunction_out();
+			break;
+		case spirv_op_typeevent:
+			out("%%%u = %s\n", opds[0], op_name);
+			break;
+		case spirv_op_typedeviceevent:
+			out("%%%u = %s\n", opds[0], op_name);
+			break;
+		case spirv_op_typereserveid:
+			out("%%%u = %s\n", opds[0], op_name);
+			break;
+		case spirv_op_typequeue:
+			out("%%%u = %s\n", opds[0], op_name);
+			break;
+		case spirv_op_typepipe:
+			out("%%%u = %s access_qualifer=%s\n", opds[0], op_name, access_qualifier_to_str(opds[1]));
+			break;
+		case spirv_op_typeforwardpointer:
+			out("%%%u = %s pointer_type=%%%u storage_class=%s\n", opds[0], op_name, opds[1], storage_class_to_str(opds[2]));
+			break;
+		case spirv_op_typepipestorage:
+			out("%%%u = %s\n", opds[0], op_name);
+			break;
+		case spirv_op_typenamedbarrier:
+			out("%%%u = %s\n", opds[0], op_name);
+			break;
+		/* types end -------------------------------------------------*/
+		/* constants start -------------------------------------------*/
+		case spirv_op_constanttrue:
+			out("%%%u = %s\n", opds[0], op_name);
+			break;
+		case spirv_op_constantfalse:
+			out("%%%u = %s\n", opds[0], op_name);
+			break;
+		case spirv_op_constant:
+			op_constant_out();
+			break;
+		case spirv_op_constantcomposite:
+			op_constantcomposite_out();
+			break;
+		case spirv_op_constantsampler:
+			op_constantsampler_out();
+			break;
+		case spirv_op_constantnull:
+			out("%%%u = %s type=%%%u\n", opds[1], op_name ,opds[0]);
+			break;
+		case spirv_op_specconstanttrue:
+			out("%%%u = %s type=%%%u\n", opds[1], op_name, opds[0]);
+			break;
+		case spirv_op_specconstantfalse:
+			out("%%%u = %s type=%%%u\n", opds[1], op_name, opds[0]);
+			break;
+		case spirv_op_specconstant:
+			op_specconstant_out();
+			break;
+		case spirv_op_specconstantcomposite:
+			op_specconstantcomposite_out();
+			break;
+		case spirv_op_specconstantop:
+			op_specconstantop_out();
+			break;
+		/* constants end ---------------------------------------------*/
+		case spirv_op_variable:
+			op_variable_out(false);
+			break;
+		case spirv_op_undef:
+			out("%%%u = %s type=%%%u\n", opds[1], op_name, opds[0]);
+			break;
+		default:
+			section_breaking_op_num = op_num;
+			break;
+		}
+
+		if (section_breaking_op_num != spirv_op_invalid)
+			break;
+
+		r = read_whole_op();
+		if (r != OK)
+			break;
+	}
+	if (section_breaking_op_num == spirv_op_invalid) {
+		out("// section end: non function declarations\n");
+		exit(0);
+	}
+	out("// section end: non function declarations, breaking opcode=%s(%u)\n", ops_name[section_breaking_op_num], section_breaking_op_num);
+}
+
+#define flag(val, str) \
+if ((w & val) != 0) { \
+	if (!first) \
+		strcat(b, "|"); \
+	strcat(b, str); \
+	first = false; \
+}
+static u8 *function_control_to_str(u32 w)
+{
+	static u8 b[256];
+	bool first;
+
+	if (w == 0)
+		return "none";
+
+	first = true;
+	b[0] = 0;	
+
+	flag(0x00000001, "inline")
+	flag(0x00000002, "dont_inline")
+	flag(0x00000004, "pure")
+	flag(0x00000008, "const")
+	flag(0xfffffff0, "unknown_function_control_flag(s)")
+}
+#undef flag
+
+static u8 *scope_to_str(u32 w)
+{
+	switch (w) {
+	case 0:
+		return "cross_device";
+	case 1:
+		return "device";
+	case 2:
+		return "workgroup";
+	case 3:
+		return "subgroup";
+	case 4:
+		return "invocation";
+	case 5:
+		return "queue_family";
+	default:
+		return "unknown_scope_code";
+	}
+}
+
+#define flag_no_opd(val, str) \
+if ((opds[i] & val) != 0) { \
+	if (!first) \
+		out("|"); \
+	out(str); \
+	first = false; \
+}
+#define TMP_SZ 128
+/* return the idx of the next mem opd */
+static u16 mem_opd_out(u16 i)
+{
+	u8 tmp[TMP_SZ];
+	u16 additional_opds;
+	bool first;
+
+	additional_opds = i + 1;
+
+	if (opds[i] == 0) {
+		out("none");
+		return i + 1;
+	}
+
+	first = true;
+
+	/* order matters */
+	flag_no_opd(0x00000001, "volatile")
+
+	if ((opds[i] & 0x00000002) != 0) {
+		if (!first)
+			out("|");
+		out("aligned");
+		first = false;
+		snprintf(tmp, TMP_SZ, "(%u)", opds[additional_opds]);
+		out(tmp);
+		++additional_opds;
+	}
+
+	flag_no_opd(0x00000004, "non_temporal")
+
+	if ((opds[i] & 0x00000008) != 0) {
+		if (!first)
+			out("|");
+		out("make_pointer_available");
+		first = false;
+		snprintf(tmp, TMP_SZ, "(%s)", scope_to_str(opds[additional_opds]));
+		out(tmp);
+		++additional_opds;
+	}
+
+	if ((opds[i] & 0x00000010) != 0) {
+		if (!first)
+			out("|");
+		out("make_pointer_visible");
+		first = false;
+		snprintf(tmp, TMP_SZ, "(%s)", scope_to_str(opds[additional_opds]));
+		out(tmp);
+		++additional_opds;
+	}
+
+	flag_no_opd(0x00000010, "non_private_pointer")
+
+	flag_no_opd(0xffffffe0, "unknown_memory_operand_flag(s)-->consider the following instruction operands as corrupted")
+
+	return additional_opds;
+}
+#undef flag_no_opd
+#undef TMP_SZ
+
+/* will output the mem opds till the end of the instruction */
+static void mem_opds_out(u16 i)
+{
+	u16 mem_opd_idx;
+
+	mem_opd_idx = 0;
+	out(" memory_operands[%u]=", mem_opd_idx);
+	loop {
+		i = mem_opd_out(i);
+		++mem_opd_idx;
+
+		if (i > opds_last)
+			break;
+	}
+}
+
+static void op_load_out_depth(void)
+{
+	out_depth("%%%u = %s type=%%%u pointer=%%%u", opds[1], op_name, opds[0], opds[2]);
+	if (op_ws_n > 4)
+		mem_opds_out(3);
+	out("\n");
+}
+
+static void op_vectorshuffle_out_depth(void)
+{
+	u16 i;
+
+	out_depth("%%%u = %s type=%%%u vector_0=%%%u vector_1=%%%u", opds[1], op_name, opds[0], opds[2], opds[3]);
+
+	i = 4;
+	loop {
+		if (i > opds_last)
+			break;
+		out(" components[%u]=%u", i - 4, opds[i]);
+		++i;
+	}
+	out("\n");
+}
+
+static void op_compositeextract_out_depth(void)
+{
+	u16 i;
+
+	out_depth("%%%u = %s type=%%%u composite=%%%u", opds[1], op_name, opds[0], opds[2]);
+
+	i = 3;
+	loop {
+		if (i > opds_last)
+			break;
+		out(" indexes[%u]=%u", i - 3, opds[i]);
+		++i;
+	}
+	out("\n");
+}
+
+static void op_compositeconstruct_out_depth(void)
+{
+	u16 i;
+
+	out_depth("%%%u = %s type=%%%u", opds[1], op_name, opds[0]);
+
+	i = 2;
+	loop {
+		if (i > opds_last)
+			break;
+		out(" constituents[%u]=%%%u", i - 2, opds[i]);
+		++i;
+	}
+	out("\n");
+}
+
+static void op_store_out_depth(void)
+{
+	out_depth("%s pointer=%%%u object=%%%u", op_name, opds[0], opds[1]);
+	if (op_ws_n > 3)
+		mem_opds_out(2);
+	out("\n");
+}
+
+static void op_imagesampleimplicitlod_out_depth(void)
+{
+	u16 i;
+	out_depth("%%%u = %s type=%%%u sampled_image=%%%u coordinate=%%%u", opds[1], op_name, opds[0], opds[2], opds[3]);
+
+	/*
+	 * XXX: this instruction encoding seems serevely broken since it
+	 * seems to depend on previous instructions, or I do not
+	 * understand how to properly decode it yet
+	 */
+	i = 4;
+	loop {
+		if (i > opds_last)
+			break;
+		out(" image_operands[%u]=0x%08x", i - 4, opds[i]);
+		++i;
+	}
+	out("\n");
+}
+
+static void op_imagesampleexplicitlod_out_depth(void)
+{
+	u16 i;
+
+	out_depth("%%%u = %s type=%%%u sampled_image=%%%u coordinate=%%%u", opds[1], op_name, opds[0], opds[2], opds[3]);
+
+	/*
+	 * XXX: this instruction encoding seems serevely broken since it
+	 * seems to depend on previous instructions, or I do not
+	 * understand how to properly decode it yet
+	 */
+	i = 4;
+	loop {
+		if (i > opds_last)
+			break;
+		out(" image_operands[%u]=0x%08x", i - 4, opds[i]);
+		++i;
+	}
+	out("\n");
+}
+
+static void op_functioncall_out_depth(void)
+{
+	u16 i;
+
+	out_depth("%%%u = %s type=%%%u function=%%%u", opds[1], op_name, opds[0], opds[2]);
+
+	i = 3;
+	loop {
+		if (i > opds_last)
+			break;
+		out(" arguments[%u]=%%%u", i - 3, opds[i]);
+		++i;
+	}	
+	out("\n");
+}
+
+static void op_branchconditional_out_depth(void)
+{
+	u16 i;
+
+	out_depth("%s condition=%%%u true_label=%%%u false_label %%%u", op_name, opds[0], opds[1], opds[2]);
+
+	i = 3;
+	loop { /* 0 or 2 */
+		if (i > opds_last)
+			break;
+		out(" weights[%u]=%u", i - 3, opds[i]);
+		++i;
+	}
+	out("\n");
+}
+
+static u8 *selection_control_to_str(u32 w)
+{
+	switch (w) {
+	case 0:
+		return "none";
+	case 1:
+		return "flatten";
+	case 2:
+		return "dont_flatten";
+	default:
+		return "unknown_selection_control_code";
+	}
+}
+
+/* TODO: add a extinst disassembler */
+static void op_extinst_out_depth(void)
+{
+	u16 i;
+
+	out_depth("%%%u = %s type=%%%u set=%%%u instruction=%u", opds[1], op_name, opds[0], opds[2], opds[3]);
+	i = 4;
+	loop {
+		if (i > opds_last)
+			break;
+		out(" operands[%u]=%%%u", i - 4, opds[i]);
+		++i;
+	}
+	out("\n");
+}
+
+static void op_accesschain_out_depth(void)
+{
+	u16 i;
+
+	out_depth("%%%u = %s type=%%%u base=%%%u", opds[1], op_name, opds[0], opds[2]);
+
+	i = 3;
+	loop {
+		if (i > opds_last)
+			break;
+		out(" indexes[%u]=%%%u", i - 3, opds[i]);
+		++i;
+	}
+	out("\n");
+}
+
+static void op_compositeinsert_out_depth(void)
+{
+	u16 i;
+
+	out_depth("%%%u = %s type=%%%u object=%%%u composite=%%%u", opds[1], op_name, opds[0], opds[2], opds[3]);
+
+	i = 4;
+	loop {
+		if (i > opds_last)
+			break;
+		out(" indexes[%u]=%u", i - 4, opds[i]);
+		++i;
+	}
+	out("\n");
+}
+
+static void op_phi_out_depth(void)
+{
+	u16 i;
+
+	out_depth("%%%u = %s type=%%%u", opds[1], op_name, opds[0]);
+
+	i = 2;
+	loop {
+		if (i > opds_last)
+			break;
+		out( "variables[%u]=%%%u parents[%u]=%%%u", (i - 2)/ 2, opds[i], opds[i + 1]);
+		i += 2;
+	}
+	out("\n");
+}
+
+/*
+ * we do a bit of state tracking in order to detect function declarations
+ * once function definitions did start, which is not allowed
+ */
+static void layout_funcs_out(void)
+{
+	u16 section_breaking_op_num;
+	bool defs_section;
+	bool func_has_blk;
+
+	defs_section = false;
+
+	section_breaking_op_num = spirv_op_invalid;
+	out("\n// section start: function declarations then definitions\n");
+	loop {
+		u8 r;
+
+		switch (op_num) {
+		case spirv_op_line:
+			out_depth("%s %%%u line=%u column=%u\n", op_name, opds[0], opds[1], opds[2]);
+			break;
+		case spirv_op_function:
+			out_depth("%%%u = %s return_type=%%%u function_control=%s function_type=%%%u\n", opds[1], op_name, opds[0], function_control_to_str(opds[2]), opds[3]);
+			func_has_blk = false;
+			++depth; /* start of a blk */
+			break;
+		case spirv_op_functionparameter:
+			out_depth("%%%u = %s type=%%%u\n", opds[1], op_name, opds[0]);
+			break;
+		case spirv_op_functionend:
+			depth--; /* end of blk */
+			out_depth("%s\n", op_name);
+			if (!func_has_blk && defs_section) {
+				out_depth("error: function declaration in function definition section\n");
+				exit(1);
+			}
+			break;
+		case spirv_op_label:
+			func_has_blk = true;
+			defs_section = true;
+			out_depth("%%%u = %s\n", opds[0], op_name);
+			++depth; /* start of a blk */
+			break;
+		case spirv_op_load:
+			op_load_out_depth();
+			break;
+		case spirv_op_vectorshuffle:
+			op_vectorshuffle_out_depth();
+			break;
+		case spirv_op_compositeextract:
+			op_compositeextract_out_depth();
+			break;
+		case spirv_op_compositeconstruct:
+			op_compositeconstruct_out_depth();
+			break;
+		case spirv_op_store:
+			op_store_out_depth();
+			break;
+		case spirv_op_return:
+			out_depth("%s\n", op_name);
+			--depth; /* end of blk */
+			break;
+		case spirv_op_imagesampleimplicitlod:
+			op_imagesampleimplicitlod_out_depth();
+			break;
+		case spirv_op_fmul:
+			out_depth("%%%u = %s type=%%%u operands[0]=%%%u operands[1]=%%%u\n", opds[1], op_name, opds[0], opds[2], opds[3]);
+			break;
+		case spirv_op_variable:
+			op_variable_out(true);
+			break;
+		case spirv_op_functioncall:
+			op_functioncall_out_depth();
+			break;
+		case spirv_op_fordlessthanequal:
+			out_depth("%%%u = %s type=%%%u operands[0]=%%%u operands[1]=%%%u\n", opds[1], op_name, opds[0], opds[2], opds[3]);
+			break;
+		case spirv_op_selectionmerge:
+			out_depth("%s merge_block=%%%u section_control=%s\n", op_name, opds[0], selection_control_to_str(opds[1]));
+			break;
+		case spirv_op_branchconditional:
+			op_branchconditional_out_depth();
+			depth--; /* end of blk */
+			break;
+		case spirv_op_fdiv:
+			out_depth("%%%u = %s type=%%%u operands[0]=%%%u operands[1]=%%%u\n", opds[1], op_name, opds[0], opds[2], opds[3]);
+			break;
+		case spirv_op_branch:
+			out_depth("%s target_label=%%%u\n", op_name, opds[0]);
+			depth--; /* end of blk */
+			break;
+		case spirv_op_fadd:
+			out_depth("%%%u = %s type=%%%u operands[0]=%%%u operands[1]=%%%u\n", opds[1], op_name, opds[0], opds[2], opds[3]);
+			break;
+		case spirv_op_extinst:
+			op_extinst_out_depth();
+			break;
+		case spirv_op_returnvalue:
+			out_depth("%s value=%%%u\n", op_name, opds[0]);
+			--depth; /* end of blk */
+			break;
+		case spirv_op_accesschain:
+			op_accesschain_out_depth();
+			break;
+		case spirv_op_compositeinsert:
+			op_compositeinsert_out_depth();
+			break;
+		case spirv_op_sampledimage:
+			out_depth("%%%u = %s type=%%%u image=%%%u sampler=%%%u\n", opds[1], op_name, opds[0], opds[2], opds[3]);
+			break;
+		case spirv_op_negate:
+			out_depth("%%%u = %s type=%%%u operand=%%%u\n", opds[1], op_name, opds[0], opds[2]);
+			break;
+		case spirv_op_vectortimesscalar:
+			out_depth("%%%u = %s type=%%%u vector=%%%u scalar=%%%u\n", opds[1], op_name, opds[0], opds[2], opds[3]);
+			break;
+		case spirv_op_fsub:
+			out_depth("%%%u = %s type=%%%u operands[0]=%%%u operands[1]=%%%u\n", opds[1], op_name, opds[0], opds[2], opds[3]);
+			break;
+		case spirv_op_fordnotequal:
+			out_depth("%%%u = %s type=%%%u operands[0]=%%%u operands[1]=%%%u\n", opds[1], op_name, opds[0], opds[2], opds[3]);
+			break;
+		case spirv_op_fordgreaterthanequal:
+			out_depth("%%%u = %s type=%%%u operands[0]=%%%u operands[1]=%%%u\n", opds[1], op_name, opds[0], opds[2], opds[3]);
+			break;
+		case spirv_op_phi:
+			op_phi_out_depth();
+			break;
+		case spirv_op_fordgreaterthan:
+			out_depth("%%%u = %s type=%%%u operands[0]=%%%u operands[1]=%%%u\n", opds[1], op_name, opds[0], opds[2], opds[3]);
+			break;
+		case spirv_op_logicalor:
+			out_depth("%%%u = %s type=%%%u operands[0]=%%%u operands[1]=%%%u\n", opds[1], op_name, opds[0], opds[2], opds[3]);
+			break;
+		case spirv_op_fordequal:
+			out_depth("%%%u = %s type=%%%u operands[0]=%%%u operands[1]=%%%u\n", opds[1], op_name, opds[0], opds[2], opds[3]); 
+			break;
+		case spirv_op_select:
+			out_depth("%%%u = %s type=%%%u condition=%%%u objects[0]=%%%u objects[1]=%%%u\n", opds[1], op_name, opds[0], opds[2], opds[3], opds[4]);
+			break;
+		case spirv_op_imagesampleexplicitlod:
+			op_imagesampleexplicitlod_out_depth();
+			break;
+		case spirv_op_fordlessthan:
+			out_depth("%%%u = %s type=%%%u operands[0]=%%%u operands[1]=%%%u\n", opds[1], op_name, opds[0], opds[2], opds[3]);
+			break;
+		case spirv_op_vectortimesmatrix:
+			out_depth("%%%u = %s type=%%%u vector=%%%u matrix=%%%u\n", opds[1], op_name, opds[0], opds[2], opds[3]);
+			break;
+		case spirv_op_dot:
+			out_depth("%%%u = %s type=%%%u vectors[0]=%%%u vectors[1]=%%%u\n", opds[1], op_name, opds[0], opds[2], opds[3]);
+			break;
+		case spirv_op_fmod:
+			out_depth("%%%u = %s type=%%%u operands[0]=%%%u operands[1]=%%%u\n", opds[1], op_name, opds[0], opds[2], opds[3]);
+			break;
+		case spirv_op_logicaland:
+			out_depth("%%%u = %s type=%%%u operands[0]=%%%u operands[1]=%%%u\n", opds[1], op_name, opds[0], opds[2], opds[3]);
+			break;
+		/* TODO: MORE! */
+		default:
+			section_breaking_op_num = op_num;
+			break;
+		}
+
+		if (section_breaking_op_num != spirv_op_invalid)
+			break;
+
+		r = read_whole_op();
+		if (r != OK)
+			break;
+	}
+	if (section_breaking_op_num == spirv_op_invalid) {
+		out("// section end: function declarations then definitions\n");
+		exit(0);
+	}
+	out("// section end: function declarations and definitions, breaking opcode=%s(%u)\n", ops_name[section_breaking_op_num], section_breaking_op_num);
+}
+
+static void init_misc(void)
+{
+	u16 i;
+
+	i = 0;
+	loop {
+		if (i == 0xffff)
+			break;
+		if (ops_name[i] == 0)
+			ops_name[i] = "unkwown_opcode";
+		++i;
+	}
+
+	depth = 0;
+}
+
+int main(void)
+{
+	u8 r;
+
+	init_misc();
+
+	clearerr(stdin);
+	hdr_out();
+
+	r = read_whole_op();	
+	if (r != OK)
+		exit(0);
+	
+	/* logical layout of a spirv module ----------------------------------*/
+	layout_caps_out();
+	layout_extns_out();
+	layout_extinstimports_out();
+	layout_memorymodel_out();
+	layout_entrypoints_out();
+	layout_executionmodes_out();
+	layout_debug_out();
+	layout_annotation_out();
+	/* from here opline is allowed */
+	layout_nonfuncdecls_out();
+	layout_funcs_out();
+	/* -------------------------------------------------------------------*/
+	return 0;
+}
-- 
2.11.4.GIT