sylware / vulkan-misc (public) (License: public domain) (since 2019-10-25) (hash sha1)
misc vulkan based project
Clone URLs: https://rocketgit.com/user/sylware/vulkan-misc ssh://rocketgit@ssh.rocketgit.com/user/sylware/vulkan-misc git://git.rocketgit.com/user/sylware/vulkan-misc
master
List of commits:
Subject Hash Author Date (UTC)
braindead partial spir-v disassembler a347cbcc8284301ad6989d486373f0d2ffb9657c Sylvain BERTRAND 2019-11-16 17:48:29
Initial commit 1cd013eaa1466522d61c7b0e64d5858055fedbfa Sylvain BERTRAND 2019-10-25 18:49:16
Commit a347cbcc8284301ad6989d486373f0d2ffb9657c - braindead partial spir-v disassembler
Author: Sylvain BERTRAND
Author date (UTC): 2019-11-16 17:48
Committer name: Sylvain BERTRAND
Committer date (UTC): 2019-11-16 17:48
Parent(s): 1cd013eaa1466522d61c7b0e64d5858055fedbfa
Signing key:
Tree: 464d962dbba194ffab554f125cca6cfd97bdcbb3
File Lines added Lines deleted
spirv/dis/README 8 0
spirv/dis/dis.c 2353 0
File spirv/dis/README added (mode: 100644) (index 0000000..cc6a282)
1 braindead partial spir-v disassembler.
2
3 PROS:
4 - it is not brainf*cked hate incarnation, namely not c++
5 - it is __simple__ C
6
7 CONS:
8 - it is not RISC-V assembly
File spirv/dis/dis.c added (mode: 100644) (index 0000000..c6ee9ab)
1 #include <stdbool.h>
2 #include <stdio.h>
3 #include <stdint.h>
4 #include <stdlib.h>
5 #include <string.h>
6 /*
7 * ABBREVIATIONS:
8 * blk(s) : BLocK(S)
9 * cap(s) : CAPability(ieS)
10 * decl : DECLaration
11 * def(s) | DEFinition(S)
12 * ext : EXTented
13 * extn(s) : EXTensioN(S)
14 * hdr : HeaDeR
15 * id(s) : IDentifier(S)
16 * img(s) : IMaGe(S)
17 * inst(s) : INSTruction(S)
18 * intf(s) : INTerFace(S)
19 * lang :LANGuage
20 * lit(s): LITeral(S)
21 * litstr : LITeral STRing
22 * n : couNt
23 * num : NUMber
24 * op(s) : OPcode(S)
25 * opd(s) : OPeranD)S)
26 * src : SouRCe
27 * str : STRing
28 * sz : SiZe
29 * w(s) : Word(S)
30 */
31 #define u8 uint8_t
32 #define u16 uint16_t
33 #define u32 uint32_t
34 #define constant enum
35 #define loop for(;;)
36 #define out(fmt,...) fprintf(stdout,fmt, ##__VA_ARGS__)
37 #define out_depth(fmt,...) \
38 { \
39 u8 d; \
40 d = 0; \
41 loop { \
42 if (d == depth) \
43 break; \
44 fprintf(stdout, "\t"); \
45 ++d; \
46 } \
47 } \
48 fprintf(stdout,fmt, ##__VA_ARGS__)
49
50 constant {
51 spirv_op_undef = 1,
52 spirv_op_sourcecontinued = 2,
53 spirv_op_source = 3,
54 spirv_op_sourceextension = 4,
55 spirv_op_name = 5,
56 spirv_op_membername = 6,
57 spirv_op_string = 7,
58 spirv_op_line = 8,
59 spirv_op_extension = 10,
60 spirv_op_extinstimport = 11,
61 spirv_op_extinst = 12,
62 spirv_op_memorymodel = 14,
63 spirv_op_entrypoint = 15,
64 spirv_op_executionmode = 16,
65 spirv_op_capability = 17,
66 spirv_op_typevoid = 19,
67 spirv_op_typebool = 20,
68 spirv_op_typeint = 21,
69 spirv_op_typefloat = 22,
70 spirv_op_typevector = 23,
71 spirv_op_typematrix = 24,
72 spirv_op_typeimage = 25,
73 spirv_op_typesampler = 26,
74 spirv_op_typesampledimage = 27,
75 spirv_op_typearray = 28,
76 spirv_op_typeruntimearray = 29,
77 spirv_op_typestruct= 30,
78 spirv_op_typeopaque= 31,
79 spirv_op_typepointer= 32,
80 spirv_op_typefunction= 33,
81 spirv_op_typeevent = 34,
82 spirv_op_typedeviceevent = 35,
83 spirv_op_typereserveid = 36,
84 spirv_op_typequeue = 37,
85 spirv_op_typepipe = 38,
86 spirv_op_typeforwardpointer = 39,
87 spirv_op_constanttrue = 41,
88 spirv_op_constantfalse = 42,
89 spirv_op_constant = 43,
90 spirv_op_constantcomposite = 44,
91 spirv_op_constantsampler = 45,
92 spirv_op_constantnull = 46,
93 spirv_op_specconstanttrue = 48,
94 spirv_op_specconstantfalse = 49,
95 spirv_op_specconstant = 50,
96 spirv_op_specconstantcomposite = 51,
97 spirv_op_specconstantop = 52,
98 spirv_op_function = 54,
99 spirv_op_functionparameter = 55,
100 spirv_op_functionend = 56,
101 spirv_op_functioncall = 57,
102 spirv_op_variable = 59,
103 spirv_op_load = 61,
104 spirv_op_store = 62,
105 spirv_op_accesschain = 65,
106 spirv_op_decorate = 71,
107 spirv_op_memberdecorate = 72,
108 spirv_op_decorationgroup = 73,
109 spirv_op_groupdecorate = 74,
110 spirv_op_groupmemberdecorate = 75,
111 spirv_op_vectorshuffle = 79,
112 spirv_op_compositeconstruct = 80,
113 spirv_op_compositeextract = 81,
114 spirv_op_compositeinsert = 82,
115 spirv_op_sampledimage = 86,
116 spirv_op_imagesampleimplicitlod = 87,
117 spirv_op_imagesampleexplicitlod = 88,
118 spirv_op_negate = 127,
119 spirv_op_fadd = 129,
120 spirv_op_fsub = 131,
121 spirv_op_fmul = 133,
122 spirv_op_fdiv = 136,
123 spirv_op_fmod = 141,
124 spirv_op_vectortimesscalar = 142,
125 spirv_op_vectortimesmatrix = 144,
126 spirv_op_dot = 148,
127 spirv_op_logicalor = 166,
128 spirv_op_logicaland = 167,
129 spirv_op_select = 169,
130 spirv_op_fordequal = 180,
131 spirv_op_fordnotequal = 182,
132 spirv_op_fordlessthan = 184,
133 spirv_op_fordgreaterthan = 186,
134 spirv_op_fordlessthanequal = 188,
135 spirv_op_fordgreaterthanequal = 190,
136 spirv_op_phi = 245,
137 spirv_op_selectionmerge = 247,
138 spirv_op_label = 248,
139 spirv_op_branch = 249,
140 spirv_op_branchconditional = 250,
141 spirv_op_return = 253,
142 spirv_op_returnvalue = 254,
143 spirv_op_typepipestorage = 322,
144 spirv_op_typenamedbarrier = 327,
145 spirv_op_moduleprocessed = 330,
146 spirv_op_executionmodeid = 331,
147 spirv_op_invalid = 0x0000ffff,
148 spirv_op_max = 0x7fffffff
149 };
150
151 /* may be augmented one day to a ops database */
152 u8 *ops_name[0xffff + 1] = {
153 [spirv_op_undef] = "undef",
154 [spirv_op_sourcecontinued] = "source_continued",
155 [spirv_op_source] = "source",
156 [spirv_op_sourceextension] = "source_extension",
157 [spirv_op_name] = "name",
158 [spirv_op_membername] = "member_name",
159 [spirv_op_string] = "string",
160 [spirv_op_line] = "line",
161 [spirv_op_extension] = "extension",
162 [spirv_op_extinstimport] = "ext_inst_import",
163 [spirv_op_extinst] = "ext_inst",
164 [spirv_op_memorymodel] = "memory_model",
165 [spirv_op_entrypoint] = "entry_point",
166 [spirv_op_executionmode] = "execution_mode",
167 [spirv_op_capability] = "capability",
168 [spirv_op_typevoid] = "type_void",
169 [spirv_op_typebool] = "type_bool",
170 [spirv_op_typeint] = "type_int",
171 [spirv_op_typefloat] = "type_float",
172 [spirv_op_typevector] = "type_vector",
173 [spirv_op_typematrix] = "type_matrix",
174 [spirv_op_typeimage] = "type_image",
175 [spirv_op_typesampler] = "type_sampler",
176 [spirv_op_typesampledimage] = "type_sampled_image",
177 [spirv_op_typearray] = "type_array",
178 [spirv_op_typeruntimearray] = "type_runtime_array",
179 [spirv_op_typestruct] = "type_struct",
180 [spirv_op_typeopaque] = "type_opaque",
181 [spirv_op_typepointer] = "type_pointer",
182 [spirv_op_typefunction] = "type_function",
183 [spirv_op_typeevent] = "type_event",
184 [spirv_op_typedeviceevent] = "type_device_event",
185 [spirv_op_typereserveid] = "type_reserve_id",
186 [spirv_op_typequeue] = "type_queue",
187 [spirv_op_typepipe] = "type_pipe",
188 [spirv_op_typeforwardpointer] = "type_forward_pointer",
189 [spirv_op_constanttrue] = "constant_true",
190 [spirv_op_constantfalse] = "constant_false",
191 [spirv_op_constant] = "constant",
192 [spirv_op_constantcomposite] = "constant_composite",
193 [spirv_op_constantsampler] = "constant_sampler",
194 [spirv_op_constantnull] = "constant_null",
195 [spirv_op_specconstanttrue] = "spec_constant_true",
196 [spirv_op_specconstantfalse] = "spec_constant_false",
197 [spirv_op_specconstant] = "spec_constant",
198 [spirv_op_specconstantcomposite] = "spec_constant_composite",
199 [spirv_op_specconstantop] = "spec_constant_op",
200 [spirv_op_function] = "function",
201 [spirv_op_functionparameter] = "function_parameter",
202 [spirv_op_functionend] = "function_end",
203 [spirv_op_functioncall] = "function_call",
204 [spirv_op_variable] = "variable",
205 [spirv_op_load] = "load",
206 [spirv_op_store] = "store",
207 [spirv_op_accesschain] = "accesschain",
208 [spirv_op_decorate] = "decorate",
209 [spirv_op_memberdecorate] = "member_decorate",
210 [spirv_op_decorationgroup] = "decoration_group",
211 [spirv_op_groupdecorate] = "group_decorate",
212 [spirv_op_groupmemberdecorate] = "group_member_decorate",
213 [spirv_op_vectorshuffle] = "vector_shuffle",
214 [spirv_op_compositeconstruct] = "composite_construct",
215 [spirv_op_compositeextract] = "composite_extract",
216 [spirv_op_compositeinsert] = "composite_insert",
217 [spirv_op_sampledimage] = "sampled_image",
218 [spirv_op_imagesampleimplicitlod] = "image_sample_implicit_lod",
219 [spirv_op_imagesampleexplicitlod] = "image_sample_explicit_lod",
220 [spirv_op_negate] = "negate",
221 [spirv_op_fadd] = "fadd",
222 [spirv_op_fsub] = "fsub",
223 [spirv_op_fmul] = "fmul",
224 [spirv_op_fdiv] = "fdiv",
225 [spirv_op_fmod] = "fmod",
226 [spirv_op_vectortimesscalar] = "vector_times_scalar",
227 [spirv_op_vectortimesmatrix] = "vector_times_matrix",
228 [spirv_op_dot] = "dot",
229 [spirv_op_logicalor] = "logical_or",
230 [spirv_op_logicaland] = "logical_and",
231 [spirv_op_select] = "select",
232 [spirv_op_fordequal] = "ford_equal",
233 [spirv_op_fordnotequal] = "ford_not_equal",
234 [spirv_op_fordlessthan] = "ford_less_than",
235 [spirv_op_fordgreaterthan] = "ford_greater_than",
236 [spirv_op_fordlessthanequal] = "ford_less_than_equal",
237 [spirv_op_fordgreaterthanequal] = "ford_greater_than_equal",
238 [spirv_op_phi] = "phi",
239 [spirv_op_selectionmerge] = "selection_merge",
240 [spirv_op_label] = "label",
241 [spirv_op_branch] = "branch",
242 [spirv_op_branchconditional] = "branch_conditional",
243 [spirv_op_return] = "return",
244 [spirv_op_returnvalue] = "return_value",
245 [spirv_op_typepipestorage] = "type_pipe_storage",
246 [spirv_op_typenamedbarrier] = "type_named_barrier",
247 [spirv_op_moduleprocessed] = "module_processed",
248 [spirv_op_executionmodeid] = "execution_mode_id",
249 [spirv_op_invalid] = "invalid"
250 };
251
252 struct hdr {
253 u32 magic_num;
254 u32 version;
255 u32 generator;
256 u32 bound;
257 u32 reserved;
258 };
259
260 static u8 depth;
261
262 static u32 op;
263 #define op_ws_n (op >> 16)
264 #define op_ws_last (op_ws_n - 1) /* last op w idx */
265 #define op_ws_end op_ws_n /* idx right after the last op w */
266 #define op_num (op & 0x0000ffff)
267
268 static u32 opds[0x0000ffff + 1];
269 #define opds_n (op_ws_n - 1)
270 #define opds_last (opds_n - 1) /* last opd w idx */
271 #define opds_end opds_n /* w idx right after the last opd w*/
272
273 #define op_name ops_name[op_num]
274
275 #define OK 0
276 #define END 1
277 #define ERR 2
278 static u8 read_whole_op(void)
279 {
280 size_t read_bytes_n;
281 read_bytes_n = fread(&op, sizeof(op), 1, stdin);
282 if (read_bytes_n != sizeof(op)) {
283 if (feof(stdin) != 0)
284 return END;
285 if (ferror(stdin) != 0)
286 return ERR;
287 }
288
289 read_bytes_n = fread(opds, sizeof(u32), opds_n, stdin);
290 if (read_bytes_n != (sizeof(u32) * opds_n)) {
291 if (feof(stdin) != 0)
292 return END;
293 if (ferror(stdin) != 0)
294 return ERR;
295 }
296 }
297
298 static void hdr_out(void)
299 {
300 struct hdr hdr;
301
302 (void)fread(&hdr, sizeof(hdr), 1, stdin);
303 if (feof(stdin) != 0 || ferror(stdin) != 0)
304 exit(0);
305
306 out("// magic number = 0x%08x (%s)\n", hdr.magic_num, hdr.magic_num == 0x07230203 ? "good" : "bad");
307 out("// version = 0x%08x\n", hdr.version);
308 out("// generator = 0x%08x\n", hdr.generator);
309 out("// bound = %u\n", hdr.bound);
310 out("// reserved = 0x%08x\n", hdr.reserved);
311 }
312
313 static void layout_caps_out(void)
314 {
315 u16 section_breaking_op_num;
316
317 section_breaking_op_num = spirv_op_invalid;
318 out("\n// section start: capabilities\n");
319 loop {
320 u8 r;
321
322 if (op_num != spirv_op_capability) {
323 section_breaking_op_num = op_num;
324 break;
325 }
326
327 out("%s capability=0x%08x\n", op_name, opds[0]);
328
329 r = read_whole_op();
330 if (r != OK)
331 break;
332 }
333 if (section_breaking_op_num == spirv_op_invalid) {
334 out("// section end: capabilities\n");
335 exit(0);
336 }
337 out("// section end: capabilities, breaking opcode=%s(%u)\n", ops_name[section_breaking_op_num], section_breaking_op_num);
338 }
339
340 /* return the sz of the litstr as a n of ws */
341 static u16 litstr_out(u32 *p)
342 {
343 u8 b[4];
344 u32 *w;
345 u16 ws_n;
346
347 w = (u32*)b;
348 out("\"");
349 ws_n = 1;
350 loop {
351 u8 i;
352
353 *w = *p;
354 i = 0;
355 loop {
356 if (i == 4)
357 break;
358
359 if (b[i] == 0) { /* 0 terminating char */
360 out("\"");
361 return ws_n;
362 }
363
364 (void)fwrite(&b[i], 1, 1, stdout);
365
366 ++i;
367 }
368 ++p;
369 ++ws_n;
370 }
371 }
372
373 static void layout_extns_out(void)
374 {
375 u16 section_breaking_op_num;
376
377 section_breaking_op_num = spirv_op_invalid;
378 out("\n// section start: extensions\n");
379 loop {
380 u8 r;
381
382 if (op_num != spirv_op_extension) {
383 section_breaking_op_num = op_num;
384 break;
385 }
386
387 out("%s ", op_name);
388 (void)litstr_out(opds);
389 out("\n");
390
391 r = read_whole_op();
392 if (r != OK)
393 break;
394 }
395 if (section_breaking_op_num == spirv_op_invalid) {
396 out("// section end: extensions\n");
397 exit(0);
398 }
399 out("// section end: extensions, breaking opcoder=%s(%u)\n", ops_name[section_breaking_op_num], section_breaking_op_num);
400 }
401
402 static void layout_extinstimports_out(void)
403 {
404 u16 section_breaking_op_num;
405
406 section_breaking_op_num = spirv_op_invalid;
407 out("\n// section start: extended set of instructions imports\n");
408 loop {
409 u8 r;
410
411 if (op_num != spirv_op_extinstimport) {
412 section_breaking_op_num = op_num;
413 break;
414 }
415
416 out("%%%u = %s name=", opds[0], op_name);
417 litstr_out(opds + 1);
418 out("\n");
419
420 r = read_whole_op();
421 if (r != OK)
422 break;
423 }
424 if (section_breaking_op_num == spirv_op_invalid) {
425 out("// section end: extended set of instructions imports\n");
426 exit(0);
427 }
428 out("// section end: extended set of instructions imports, breaking opcode=%s(%u)\n", ops_name[section_breaking_op_num], section_breaking_op_num);
429 }
430
431 static u8 *addressing_model_to_str(u32 w)
432 {
433 switch (w) {
434 case 0:
435 return "logical";
436 case 1:
437 return "physical32";
438 case 2:
439 return "physical64";
440 default:
441 return "unknown_addressing_model_code";
442 }
443 }
444
445 static u8 *memory_model_to_str(u32 w)
446 {
447 switch (w) {
448 case 0:
449 return "simple";
450 case 1:
451 return "glsl450";
452 case 2:
453 return "opencl";
454 case 3:
455 return "vulkan";
456 default:
457 return "unkwnown_memory_model_code";
458 }
459 }
460
461 static void layout_memorymodel_out(void)
462 {
463 u8 r;
464
465 out("\n// the only memory model instruction, if one, should be here\n");
466 if (op_num == spirv_op_memorymodel)
467 out("%s addressing_model=%s memory_model=%s\n", op_name, addressing_model_to_str(opds[0]), memory_model_to_str(opds[1]));
468
469 r = read_whole_op();
470 if (r != OK)
471 exit(0);
472 }
473
474 static u8 *execution_model_to_str(u32 w)
475 {
476 switch (w) {
477 case 0:
478 return "vertex";
479 case 1:
480 return "tessellation_control";
481 case 2:
482 return "tessellation_evaluation";
483 case 3:
484 return "geometry";
485 case 4:
486 return "fragment";
487 case 5:
488 return "gl_compute";
489 case 6:
490 return "kernel";
491 default:
492 return "unkwown_execution_model_code";
493 }
494 }
495
496 static void layout_entrypoints_out(void)
497 {
498 u16 section_breaking_op_num;
499
500 section_breaking_op_num = spirv_op_invalid;
501 out("\n// section start: entry points\n");
502 loop {
503 u8 r;
504 u16 name_ws_n;
505 u16 intf_ws_n;
506
507 if (op_num != spirv_op_entrypoint) {
508 section_breaking_op_num = op_num;
509 break;
510 }
511
512 out("%s execution_model=%s entry_point=%%%u name=", op_name, execution_model_to_str(opds[0]), opds[1]);
513 name_ws_n = litstr_out(opds + 2);
514
515 intf_ws_n = op_ws_n - (3 + name_ws_n);
516 if (intf_ws_n != 0) {
517 u16 i;
518
519 i = 0;
520 loop {
521 if (i == intf_ws_n)
522 break;
523 out(" interfaces[%u]=%%%u", i, opds[2 + name_ws_n + i]);
524 ++i;
525 }
526 }
527 out("\n");
528
529 r = read_whole_op();
530 if (r != OK)
531 break;
532 }
533 if (section_breaking_op_num == spirv_op_invalid) {
534 out("// section end: entry points\n");
535 exit(0);
536 }
537 out("// section end: entry points, breaking opcode=%s(%u)\n", ops_name[section_breaking_op_num], section_breaking_op_num);
538 }
539
540 static void op_executionmode_out(void)
541 {
542 u16 execution_mode_lits_ws_n;
543 u16 i;
544
545 /*
546 * we don't decode the execution modes: we'll add the ones we are
547 * interested in
548 */
549 out("%s entry_point=%%%u mode=0x%08x\n", op_name, opds[0], opds[1]);
550
551 execution_mode_lits_ws_n = op_ws_n - 3;
552 if (execution_mode_lits_ws_n == 0 )
553 return;
554
555 i = 0;
556 loop {
557 if (i == execution_mode_lits_ws_n)
558 break;
559 out(" 0x%08x", opds[2 + i]);
560 ++i;
561 }
562 }
563
564 static void op_executionmodeid_out(void)
565 {
566 u16 execution_mode_ids_ws_n;
567 u16 i;
568
569 out("%s %%%u 0x%08x ", op_name, opds[0], opds[1]);
570
571 execution_mode_ids_ws_n = op_ws_n - 3;
572 if (execution_mode_ids_ws_n == 0)
573 return;
574
575 i = 0;
576 loop {
577 if (i == execution_mode_ids_ws_n)
578 break;
579 out(" %%%u", opds[2 + i]);
580 ++i;
581 }
582 }
583
584 static void layout_executionmodes_out(void)
585 {
586 u16 section_breaking_op_num;
587
588 section_breaking_op_num = spirv_op_invalid;
589 out("\n// section start: execution modes\n");
590 loop {
591 u8 r;
592
593 if (op_num == spirv_op_executionmode) {
594 op_executionmode_out();
595 } else if (op_num == spirv_op_executionmodeid) {
596 op_executionmodeid_out();
597 } else {
598 section_breaking_op_num = op_num;
599 break;
600 }
601
602 r = read_whole_op();
603 if (r != OK)
604 break;
605 }
606 if (section_breaking_op_num == spirv_op_invalid) {
607 out("// section end: execution modes\n");
608 exit(0);
609 }
610 out("// section end: execution modes, breaking opcoder=%s(%u)\n", ops_name[section_breaking_op_num], section_breaking_op_num);
611 }
612
613 static u8 *src_lang_str(u32 w)
614 {
615 switch (w) {
616 case 0:
617 return "unknown";
618 case 1:
619 return "essl";
620 case 2:
621 return "glsl";
622 case 3:
623 return "opencl_c";
624 case 4:
625 return "opencl_cpp";
626 case 5:
627 return "hlsl";
628 default:
629 return "implicit_unknow";
630 }
631 }
632
633 static void op_sourcecontinued_out(void)
634 {
635 out("%s ", op_name);
636 (void)litstr_out(opds);
637 out("\n");
638 }
639
640 static void op_source_out(void)
641 {
642 out("5s %s 0x%08x", op_name, src_lang_str(opds[0]), opds[1]);
643 if (op_ws_n > 3) {
644 out(" %%%u", opds[2]);
645
646 if (op_ws_n > 4) {
647 out(" ");
648 (void)litstr_out(opds + 3);
649 }
650 }
651 out("\n");
652 }
653
654 static void op_sourceextension_out(void)
655 {
656 out("%s ", op_name);
657 (void)litstr_out(opds);
658 out("\n");
659 }
660
661 static void op_string_out(void)
662 {
663 out("%s %%%u", op_name, opds[0]);
664 (void)litstr_out(opds + 1);
665 out("\n");
666 }
667
668 static bool layout_debug_section_0(void)
669 {
670 u16 section_breaking_op_num;
671
672 section_breaking_op_num = spirv_op_invalid;
673 out("\n// debug first subsection start\n");
674 loop {
675 u8 r;
676
677 if (op_num == spirv_op_sourcecontinued) {
678 op_sourcecontinued_out();
679 } else if (op_num == spirv_op_source) {
680 op_source_out();
681 } else if (op_num == spirv_op_sourceextension) {
682 op_sourceextension_out();
683 } else if (op_num == spirv_op_string) {
684 op_string_out();
685 } else {
686 section_breaking_op_num = op_num;
687 break;
688 }
689
690 r = read_whole_op();
691 if (r != OK)
692 break;
693 }
694 if (section_breaking_op_num == spirv_op_invalid) {
695 out("// debug first subsection end\n");
696 return true;
697 }
698 out("// debug first subsection end, breaking opcode=%s(%u)\n", ops_name[section_breaking_op_num], section_breaking_op_num);
699 return false;
700 }
701
702 static void op_name_out(void)
703 {
704 out("%s target=%%%u name=", op_name, opds[0]);
705 (void)litstr_out(opds + 1);
706 out("\n");
707 }
708
709 static void op_membername_out(void)
710 {
711 out("%s type=%%%u member=%u name=", op_name, opds[0], opds[1]);
712 (void)litstr_out(opds + 2);
713 out("\n");
714 }
715
716 static bool layout_debug_section_1(void)
717 {
718 u16 section_breaking_op_num;
719
720 section_breaking_op_num = spirv_op_invalid;
721 out("\n// debug second subsection start\n");
722 loop {
723 u8 r;
724
725 if (op_num == spirv_op_name) {
726 op_name_out();
727 } else if (op_num == spirv_op_membername) {
728 op_membername_out();
729 } else {
730 section_breaking_op_num = op_num;
731 break;
732 }
733
734 r = read_whole_op();
735 if (r != OK)
736 break;
737 }
738 if (section_breaking_op_num == spirv_op_invalid) {
739 out("// debug second subsection end\n");
740 return true;
741 }
742 out("// debug second subsection end, breaking opcode=%s(%u)\n", ops_name[section_breaking_op_num], section_breaking_op_num);
743 return false;
744 }
745
746 static bool layout_debug_section_2(void)
747 {
748 u16 section_breaking_op_num;
749
750 section_breaking_op_num = spirv_op_invalid;
751 out("\n// debug third subsection start\n");
752 loop {
753 u8 r;
754
755 if (op_num != spirv_op_moduleprocessed) {
756 section_breaking_op_num = op_num;
757 break;
758 }
759 out("%s ", op_name);
760 (void)litstr_out(opds);
761 out("\n");
762
763 r = read_whole_op();
764 if (r != OK)
765 break;
766 }
767 if (section_breaking_op_num == spirv_op_invalid) {
768 out("// debug third subsection end\n");
769 return true;
770 }
771 out("// debug third subsection end, breaking opcode=%s(%u)\n", ops_name[section_breaking_op_num], section_breaking_op_num);
772 return false;
773 }
774
775 static void layout_debug_out(void)
776 {
777 bool do_exit;
778
779 out("\n// section start: debug\n");
780
781 do_exit = layout_debug_section_0();
782 if (do_exit)
783 goto exit;
784
785 do_exit = layout_debug_section_1();
786 if (do_exit)
787 goto exit;
788
789 do_exit = layout_debug_section_2();
790
791 exit:
792 out("// section end: debug\n");
793 if (do_exit)
794 exit(0);
795 }
796
797 static void decoration_builtin_out(u16 i)
798 {
799 u8 *str;
800
801 switch (opds[i + 1]) {
802 case 0:
803 str = "position";
804 break;
805 case 1:
806 str = "point_size";
807 break;
808 case 3:
809 str = "clip_distance";
810 break;
811 case 4:
812 str = "cull_distance";
813 break;
814 case 5:
815 str = "vertex_id";
816 break;
817 case 6:
818 str = "instance_id";
819 break;
820 case 7:
821 str = "primitive_id";
822 break;
823 case 8:
824 str = "invocation_id";
825 break;
826 case 9:
827 str = "layer";
828 break;
829 case 10:
830 str = "viewport_idx";
831 break;
832 case 11:
833 str = "tess_level_outer";
834 break;
835 case 12:
836 str = "tess_level_inner";
837 break;
838 case 13:
839 str = "tess_coord";
840 break;
841 case 14:
842 str = "patch_vertices";
843 break;
844 case 15:
845 str = "frag_coord";
846 break;
847 case 16:
848 str = "point_coord";
849 break;
850 case 17:
851 str = "front_facing";
852 break;
853 case 18:
854 str = "sample_id";
855 break;
856 case 19:
857 str = "sample_position";
858 break;
859 case 20:
860 str = "sample_mask";
861 break;
862 case 22:
863 str = "frag_depth";
864 break;
865 case 23:
866 str = "helper_invocation";
867 break;
868 case 24:
869 str = "num_workgroups";
870 break;
871 case 25:
872 str = "workgroup_size";
873 break;
874 case 26:
875 str = "workgroup_id";
876 break;
877 case 27:
878 str = "local_invocation_id";
879 break;
880 case 28:
881 str = "global_invocation_id";
882 break;
883 case 29:
884 str = "global_invocation_idx";
885 break;
886 case 30:
887 str = "work_dim";
888 break;
889 case 31:
890 str = "global_size";
891 break;
892 case 32:
893 str = "enqueue_workgroup_size";
894 break;
895 case 33:
896 str = "global_offset";
897 break;
898 case 34:
899 str = "global_linear_id";
900 break;
901 case 36:
902 str = "subgroup_size";
903 break;
904 case 37:
905 str = "subgroup_max_size";
906 break;
907 case 38:
908 str = "num_subgroups";
909 break;
910 case 39:
911 str = "num_enqueued_subgroups";
912 break;
913 case 40:
914 str = "subgroup_id";
915 break;
916 case 41:
917 str = "subgroup_local_invocation_id";
918 break;
919 case 42:
920 str = "vertex_index";
921 break;
922 case 43:
923 str = "instance_index";
924 break;
925 case 4424:
926 str = "base_vertex";
927 break;
928 case 4425:
929 str = "base_instance";
930 break;
931 case 4426:
932 str = "draw_index";
933 break;
934 case 4438:
935 str = "device_index";
936 break;
937 case 4440:
938 str = "view_index";
939 break;
940 default:
941 str = "unknown_builtin_code";
942 break;
943 }
944
945 out(" builtin %s", str);
946 }
947
948 static u8 *scopeid_to_str(u32 w)
949 {
950 switch (w) {
951 case 0:
952 return "cross_device";
953 case 1:
954 return "device";
955 case 2:
956 return "workgroup";
957 case 3:
958 return "subgroup";
959 case 4:
960 return "invocation";
961 case 5:
962 return "queue_family";
963 default:
964 return "unknow_scope_id_code";
965 }
966 }
967
968 static u8 *fproundingmode_to_str(u32 w)
969 {
970 switch (w) {
971 case 0:
972 return "rte";
973 case 1:
974 return "rtz";
975 case 2:
976 return "rtp";
977 case 3:
978 return "rtn";
979 default:
980 return "unknown_fp_rounding_mode_code";
981 }
982 }
983
984 #define flag(val, str) \
985 if ((w & val) != 0) { \
986 if (!first) \
987 strcat(b, "|"); \
988 strcat(b, str); \
989 first = false; \
990 }
991 static u8 *fpfastmathmode_to_str(u32 w)
992 {
993 static u8 b[256];
994 bool first;
995
996 if (w == 0)
997 return "none";
998
999 b[0] = 0;
1000 first = true;
1001
1002 flag(0x00000001, "not_nan")
1003 flag(0x00000002, "not_inf")
1004 flag(0x00000004, "nsz")
1005 flag(0x00000008, "allow_recip")
1006 flag(0x00000010, "fast")
1007 flag(0xffffffe0, "unknown_fp_fast_math_mode_code")
1008 }
1009 #undef flag
1010
1011 static u8 *funcparamattr_to_str(u32 w)
1012 {
1013 switch (w) {
1014 case 0:
1015 return "z_ext";
1016 case 1:
1017 return "s_ext";
1018 case 2:
1019 return "by_val";
1020 case 3:
1021 return "s_ret";
1022 case 4:
1023 return "no_alias";
1024 case 5:
1025 return "no_capture";
1026 case 6:
1027 return "no_write";
1028 case 7:
1029 return "no_read_write";
1030 defaut:
1031 return "unknown_func_param_attr_code";
1032 }
1033 }
1034
1035
1036 /* i is the start idx in the opds array of the decoration specifications */
1037 static void decoration_out(u16 i)
1038 {
1039 switch (opds[i]) {
1040 case 0:
1041 out(" relaxedr_|precision");
1042 break;
1043 case 1:
1044 out(" spec_id");
1045 break;
1046 case 2:
1047 out(" block");
1048 break;
1049 case 3:
1050 out(" buffer_block");
1051 break;
1052 case 4:
1053 out(" row_major");
1054 break;
1055 case 5:
1056 out(" col_major");
1057 break;
1058 case 6:
1059 out(" array_stride %u", opds[i + 1]);
1060 break;
1061 case 7:
1062 out(" matrix_stride %u", opds[i + 1]);
1063 break;
1064 case 8:
1065 out(" glsl_shared");
1066 break;
1067 case 9:
1068 out(" glsl_packed");
1069 break;
1070 case 10:
1071 out(" c_packed");
1072 break;
1073 case 11:
1074 decoration_builtin_out(i);
1075 break;
1076 case 13:
1077 out(" no_perspective");
1078 break;
1079 case 14:
1080 out(" flat");
1081 break;
1082 case 15:
1083 out(" patch");
1084 break;
1085 case 16:
1086 out(" centroid");
1087 break;
1088 case 17:
1089 out(" sample");
1090 break;
1091 case 18:
1092 out(" invariant");
1093 break;
1094 case 19:
1095 out(" restrict");
1096 break;
1097 case 20:
1098 out(" aliased");
1099 break;
1100 case 21:
1101 out(" volatile");
1102 break;
1103 case 22:
1104 out(" constant");
1105 break;
1106 case 23:
1107 out(" coherent");
1108 break;
1109 case 24:
1110 out(" nonwritable");
1111 break;
1112 case 25:
1113 out(" nonreadable");
1114 break;
1115 case 26:
1116 out(" uniform");
1117 break;
1118 case 27:
1119 out(" uniform_id %s", scopeid_to_str(opds[i + 1]));
1120 break;
1121 case 28:
1122 out(" saturated_conversion");
1123 break;
1124 case 29:
1125 out(" stream %u", opds[i + 1]);
1126 break;
1127 case 30:
1128 out(" location %u", opds[i + 1]);
1129 break;
1130 case 31:
1131 out(" component %u", opds[i + 1]);
1132 break;
1133 case 32:
1134 out(" index %u", opds[i + 1]);
1135 break;
1136 case 33:
1137 out(" binding %u", opds[i + 1]);
1138 break;
1139 case 34:
1140 out(" descriptor_set %u", opds[i + 1]);
1141 break;
1142 case 35:
1143 out(" offset %u", opds[i + 1]);
1144 break;
1145 case 36:
1146 out(" xfb_buffer %u", opds[i + 1]);
1147 break;
1148 case 37:
1149 out(" xfb_stride %u", opds[i + 1]);
1150 break;
1151 case 38:
1152 out(" func_param_attr %s", funcparamattr_to_str(opds[i + 1]));
1153 break;
1154 case 39:
1155 out(" fp_rounding_mode %s", fproundingmode_to_str(opds[i + 1]));
1156 break;
1157 case 40:
1158 out(" fp_fast_math_mode %s", fpfastmathmode_to_str(opds[i + 1]));
1159 break;
1160 case 41:
1161 out(" linkage_attributes (...)");
1162 break;
1163 case 42:
1164 out(" no_contraction");
1165 break;
1166 case 43:
1167 out(" input_attachment_index %u", opds[i + 1]);
1168 break;
1169 case 44:
1170 out(" alignment %u", opds[i + 1]);
1171 break;
1172
1173 case 45:
1174 out(" max_byte_offset %u", opds[i + 1]);
1175 break;
1176 case 46:
1177 out(" alignment_id %%%u", opds[i + 1]);
1178 break;
1179 case 47:
1180 out(" max_byte_offset_id %%%u", opds[i + 1]);
1181 break;
1182 default:
1183 out(" decoration_not_handled(%u operands)", op_ws_n - i - 1);
1184 break;
1185 }
1186 }
1187
1188 static void op_decorate_out(void)
1189 {
1190 out("%s %%%u", op_name, opds[0]);
1191 decoration_out(1);
1192 out("\n");
1193 }
1194
1195 static void op_memberdecorate_out(void)
1196 {
1197 out("%s %%%u %u", op_name, opds[0], opds[1]);
1198 decoration_out(2);
1199 out("\n");
1200 }
1201
1202 static void layout_annotation_out(void)
1203 {
1204 u16 section_breaking_op_num;
1205
1206 section_breaking_op_num = spirv_op_invalid;
1207 out("\n// section start: annotations\n");
1208 loop {
1209 u8 r;
1210
1211 if (op_num == spirv_op_decorate) {
1212 op_decorate_out();
1213 } else if (op_num == spirv_op_memberdecorate) {
1214 op_memberdecorate_out();
1215 } else if (op_num == spirv_op_groupdecorate) {
1216 out("%s(deprecated, SKIPPING)\n", op_name);
1217 } else if (op_num == spirv_op_groupmemberdecorate) {
1218 out("%s(deprecated, SKIPPING)\n", op_name);
1219 } else if (op_num == spirv_op_decorationgroup) {
1220 out("%s(deprecated, SKIPPING)\n", op_name);
1221 } else {
1222 section_breaking_op_num = op_num;
1223 break;
1224 }
1225
1226 r = read_whole_op();
1227 if (r != OK)
1228 break;
1229 }
1230 if (section_breaking_op_num == spirv_op_invalid) {
1231 out("// section end: annotations\n");
1232 exit(0);
1233 }
1234 out("// section end: annotations, breaking opcode=%s(%u)\n", ops_name[section_breaking_op_num], section_breaking_op_num);
1235 }
1236
1237 static u8 *signedness_to_str(u32 w)
1238 {
1239 switch (w) {
1240 case 0:
1241 return "unsigned";
1242 case 1:
1243 return "signed";
1244 default:
1245 return "unknown_signedness_code";
1246 }
1247 }
1248
1249 static u8 *dim_to_str(u32 w)
1250 {
1251 switch (w) {
1252 case 0:
1253 return "1d";
1254 case 1:
1255 return "2d";
1256 case 2:
1257 return "3d";
1258 case 3:
1259 return "cube";
1260 case 4:
1261 return "rect";
1262 case 5:
1263 return "buffer";
1264 case 6:
1265 return "subpass_data";
1266 default:
1267 return "unknown_dimension_code";
1268 }
1269 }
1270
1271 static u8 *depth_to_str(u32 w)
1272 {
1273 switch (w) {
1274 case 0:
1275 return "no_depth_image";
1276 case 1:
1277 return "depth_image";
1278 case 2:
1279 return "no_depth_information";
1280 default:
1281 return "unknown_depth_code";
1282 }
1283 }
1284
1285 static u8 *arrayed_to_str(u32 w)
1286 {
1287 switch (w) {
1288 case 0:
1289 return "non_arrayed_content";
1290 case 1:
1291 return "arrayed_content";
1292 default:
1293 return "unknown_arrayed_code";
1294 }
1295 }
1296
1297 static u8 *multisample_to_str(u32 w)
1298 {
1299 switch (w) {
1300 case 0:
1301 return "single_sampled";
1302 case 1:
1303 return "multi_sampled";
1304 default:
1305 return "unknown_multisample_code";
1306 }
1307 }
1308
1309 static u8 *sampled_to_str(u32 w)
1310 {
1311 switch (w) {
1312 case 0:
1313 return "runtime_known";
1314 case 1:
1315 return "sampler";
1316 case 2:
1317 return "no_sampler";
1318 default:
1319 return "unknown_sampled_code";
1320 }
1321 }
1322
1323 static u8 *img_fmt_to_str(u32 w)
1324 {
1325 switch (w) {
1326 case 0:
1327 return "unkown";
1328 case 1:
1329 return "rgba32f";
1330 case 2:
1331 return "rgba16f";
1332 case 3:
1333 return "r32f";
1334 case 4:
1335 return "rgba8";
1336 case 5:
1337 return "rgba8snorm";
1338 case 6:
1339 return "rg32f";
1340 case 7:
1341 return "rg16f";
1342 case 8:
1343 return "r11g11b10f";
1344 case 9:
1345 return "r16f";
1346 case 10:
1347 return "rgba16";
1348 case 11:
1349 return "rgb10a2";
1350 case 12:
1351 return "rg16";
1352 case 13:
1353 return "rg8";
1354 case 14:
1355 return "r16";
1356 case 15:
1357 return "r8";
1358 case 16:
1359 return "rgba16snorm";
1360 case 17:
1361 return "rg16snorm";
1362 case 18:
1363 return "rg8snorm";
1364 case 19:
1365 return "r16snorm";
1366 case 20:
1367 return "r8snorm";
1368 case 21:
1369 return "rgba32i";
1370 case 22:
1371 return "rgba16i";
1372 case 23:
1373 return "rgba8i";
1374 case 24:
1375 return "r32i";
1376 case 25:
1377 return "rg32i";
1378 case 26:
1379 return "rg16i";
1380 case 27:
1381 return "rg8i";
1382 case 28:
1383 return "r16i";
1384 case 29:
1385 return "r8i";
1386 case 30:
1387 return "rgba32ui";
1388 case 31:
1389 return "rgba16ui";
1390 case 32:
1391 return "rgba8ui";
1392 case 33:
1393 return "r32ui";
1394 case 34:
1395 return "rgb10a2ui";
1396 case 35:
1397 return "rg32ui";
1398 case 36:
1399 return "rg16ui";
1400 case 37:
1401 return "rg8ui";
1402 case 38:
1403 return "r16ui";
1404 case 39:
1405 return "r8ui";
1406 default:
1407 return "unknown_image_format_code";
1408 }
1409 }
1410
1411 static u8 *access_qualifier_to_str(u32 w)
1412 {
1413 switch (w) {
1414 case 0:
1415 return "read_only";
1416 case 1:
1417 return "write_only";
1418 case 2:
1419 return "read_write";
1420 default:
1421 return "unknown_access_qualifier_code";
1422 }
1423 }
1424
1425 static void op_typeimage_out(void)
1426 {
1427 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]));
1428 if (op_ws_n > 9) /* have access qualifier */
1429 out(" access_qualifier=%s\n", access_qualifier_to_str(opds[8]));
1430 else
1431 out("\n");
1432 }
1433
1434 static void op_typestruct_out(void)
1435 {
1436 u16 i;
1437
1438 out("%%%u = %s", opds[0], op_name);
1439 i = 1;
1440 loop {
1441 if (i > opds_last)
1442 break;
1443 out(" member_%u=%%%u", i - 1, opds[i]);
1444 ++i;
1445 }
1446 out("\n");
1447 }
1448
1449 static u8 *storage_class_to_str(u32 w)
1450 {
1451 switch (w) {
1452 case 0:
1453 return "uniform_constant";
1454 case 1:
1455 return "input";
1456 case 2:
1457 return "uniform";
1458 case 3:
1459 return "output";
1460 case 4:
1461 return "workgroup";
1462 case 5:
1463 return "cross_workgroup";
1464 case 6:
1465 return "private";
1466 case 7:
1467 return "function";
1468 case 8:
1469 return "generic";
1470 case 9:
1471 return "push_constant";
1472 case 10:
1473 return "atomic_counter";
1474 case 11:
1475 return "image";
1476 case 12:
1477 return "storage_buffer";
1478 case 5349:
1479 return "physical_storage_buffer";
1480 default:
1481 return "unknown_storage_class_code";
1482 }
1483 }
1484
1485 static void op_typefunction_out(void)
1486 {
1487 u16 i;
1488
1489 out("%%%u = %s return_type=%%%u", opds[0], op_name, opds[1]);
1490
1491 i = 2;
1492 loop {
1493 if (i > opds_last)
1494 break;
1495 out(" parameters[%u]=%%%u", i - 2, opds[i]);
1496 ++i;
1497 }
1498 out("\n");
1499 }
1500
1501 static void op_constant_out(void)
1502 {
1503 u16 i;
1504
1505 out("%%%u = %s type=%%%u", opds[1], op_name, opds[0]);
1506
1507 i = 2;
1508 loop {
1509 if (i > opds_last)
1510 break;
1511 out(" values[%u]=0x%08x", i - 2, opds[i]);
1512 ++i;
1513 }
1514 out("\n");
1515 }
1516
1517 static void op_specconstant_out(void)
1518 {
1519 u16 i;
1520
1521 out("%%%u = %s type=%%%u", opds[1], op_name, opds[0]);
1522
1523 i = 2;
1524 loop {
1525 if (i > opds_last)
1526 break;
1527 out(" values[%u]=0x%08x", i - 2, opds[i]);
1528 ++i;
1529 }
1530 out("\n");
1531 }
1532
1533 static void op_constantcomposite_out(void)
1534 {
1535 u16 i;
1536
1537 out("%%%u = %s type=%%%u", opds[1], op_name, opds[0]);
1538
1539 i = 2;
1540 loop {
1541 if (i > opds_last)
1542 break;
1543 out(" constituents[%u]=%%%u", i - 2, opds[i]);
1544 ++i;
1545 }
1546 out("\n");
1547 }
1548
1549 static void op_specconstantcomposite_out(void)
1550 {
1551 u16 i;
1552
1553 out("%%%u = %s type=%%%u", opds[1], op_name, opds[0]);
1554
1555 i = 2;
1556 loop {
1557 if (i > opds_last)
1558 break;
1559 out(" constituents[%u]=%%%u", i - 2, opds[i]);
1560 ++i;
1561 }
1562 out("\n");
1563 }
1564
1565 static u8 *sampler_addressing_mode_to_str(u32 w)
1566 {
1567 switch (w) {
1568 case 0:
1569 return "none";
1570 case 1:
1571 return "clamp_to_edge";
1572 case 2:
1573 return "clamp";
1574 case 3:
1575 return "repeat";
1576 case 4:
1577 return "repeat_mirrored";
1578 default:
1579 return "unknown_sampler_addressing_mode_code";
1580 }
1581 }
1582
1583 static u8 *sampler_filter_mode_to_str(u32 w)
1584 {
1585 switch (w) {
1586 case 0:
1587 return "nearest";
1588 case 1:
1589 return "linear";
1590 default:
1591 return "unknown_sampler_filter_mode_code";
1592 }
1593 }
1594
1595 static void op_constantsampler_out(void)
1596 {
1597 u8 *param;
1598
1599 switch (opds[3]) {
1600 case 0:
1601 param = "non_normalized";
1602 break;
1603 case 1:
1604 param = "normalized";
1605 break;
1606 default:
1607 param = "unknown_param_code";
1608 break;
1609 }
1610 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]));
1611 }
1612
1613 static void op_specconstantop_out(void)
1614 {
1615 u16 i;
1616
1617 out("%%%u = %s type=%%%u opcode=%u", opds[1], op_name, opds[0], opds[2]);
1618
1619 i = 3;
1620 loop {
1621 if (i > opds_last)
1622 break;
1623 out(" operands[%u]=0x%08x", i - 3, opds[i]);
1624 ++i;
1625 }
1626 out("\n");
1627 }
1628
1629 static void op_variable_out(bool depth)
1630 {
1631 if (depth) {
1632 out_depth("%%%u = %s pointer_type=%%%u storage_class=%s", opds[1], op_name, opds[0], storage_class_to_str(opds[2]));
1633 } else {
1634 out("%%%u = %s pointer_type=%%%u storage_class=%s", opds[1], op_name, opds[0], storage_class_to_str(opds[2]));
1635 }
1636 if (op_ws_n > 4)
1637 out(" initializer=%%%u", opds[3]);
1638 out("\n");
1639 }
1640
1641 static void layout_nonfuncdecls_out(void)
1642 {
1643 u16 section_breaking_op_num;
1644
1645 section_breaking_op_num = spirv_op_invalid;
1646 out("\n// section start: non function declarations\n");
1647 loop {
1648 u8 r;
1649
1650 switch (op_num) {
1651 case spirv_op_line:
1652 out("%s %%%u line=%u column=%u\n", op_name, opds[0], opds[1], opds[2]);
1653 break;
1654 /* types start -----------------------------------------------*/
1655 case spirv_op_typevoid:
1656 out("%%%u = %s\n", opds[0], op_name);
1657 break;
1658 case spirv_op_typebool:
1659 out("%%%u = %s\n", opds[0], op_name);
1660 break;
1661 case spirv_op_typeint:
1662 out("%%%u = %s width=%u signedness=%s\n", opds[0], op_name, opds[1], signedness_to_str(opds[2]));
1663 break;
1664 case spirv_op_typefloat:
1665 out("%%%u = %s width=%u\n", opds[0], op_name, opds[1]);
1666 break;
1667 case spirv_op_typevector:
1668 out("%%%u = %s component_type=%%%u component_count=%u\n", opds[0], op_name, opds[1], opds[2]);
1669 break;
1670 case spirv_op_typematrix:
1671 out("%%%u = %s column_type=%%%u column_count=%u\n", opds[0], op_name, opds[1], opds[2]);
1672 break;
1673 case spirv_op_typeimage:
1674 op_typeimage_out();
1675 break;
1676 case spirv_op_typesampler:
1677 out("%%%u = %s\n", opds[0], op_name);
1678 break;
1679 case spirv_op_typesampledimage:
1680 out("%%%u = %s image_type=%%%u\n", opds[0], op_name, opds[1]);
1681 break;
1682 case spirv_op_typearray:
1683 out("%%%u = %s element_type=%%%u length=%%%u\n", opds[0], op_name, opds[1]);
1684 break;
1685 case spirv_op_typeruntimearray:
1686 out("%%%u = %s element_type=%%%u\n", opds[0], op_name, opds[1]);
1687 break;
1688 case spirv_op_typestruct:
1689 op_typestruct_out();
1690 break;
1691 case spirv_op_typeopaque:
1692 out("%%%u = %s ", opds[0], op_name);
1693 (void)litstr_out(opds + 1);
1694 out("\n");
1695 break;
1696 case spirv_op_typepointer:
1697 out("%%%u = %s storage_class=%s type=%%%u\n", opds[0], op_name, storage_class_to_str(opds[1]), opds[2]);
1698 break;
1699 case spirv_op_typefunction:
1700 op_typefunction_out();
1701 break;
1702 case spirv_op_typeevent:
1703 out("%%%u = %s\n", opds[0], op_name);
1704 break;
1705 case spirv_op_typedeviceevent:
1706 out("%%%u = %s\n", opds[0], op_name);
1707 break;
1708 case spirv_op_typereserveid:
1709 out("%%%u = %s\n", opds[0], op_name);
1710 break;
1711 case spirv_op_typequeue:
1712 out("%%%u = %s\n", opds[0], op_name);
1713 break;
1714 case spirv_op_typepipe:
1715 out("%%%u = %s access_qualifer=%s\n", opds[0], op_name, access_qualifier_to_str(opds[1]));
1716 break;
1717 case spirv_op_typeforwardpointer:
1718 out("%%%u = %s pointer_type=%%%u storage_class=%s\n", opds[0], op_name, opds[1], storage_class_to_str(opds[2]));
1719 break;
1720 case spirv_op_typepipestorage:
1721 out("%%%u = %s\n", opds[0], op_name);
1722 break;
1723 case spirv_op_typenamedbarrier:
1724 out("%%%u = %s\n", opds[0], op_name);
1725 break;
1726 /* types end -------------------------------------------------*/
1727 /* constants start -------------------------------------------*/
1728 case spirv_op_constanttrue:
1729 out("%%%u = %s\n", opds[0], op_name);
1730 break;
1731 case spirv_op_constantfalse:
1732 out("%%%u = %s\n", opds[0], op_name);
1733 break;
1734 case spirv_op_constant:
1735 op_constant_out();
1736 break;
1737 case spirv_op_constantcomposite:
1738 op_constantcomposite_out();
1739 break;
1740 case spirv_op_constantsampler:
1741 op_constantsampler_out();
1742 break;
1743 case spirv_op_constantnull:
1744 out("%%%u = %s type=%%%u\n", opds[1], op_name ,opds[0]);
1745 break;
1746 case spirv_op_specconstanttrue:
1747 out("%%%u = %s type=%%%u\n", opds[1], op_name, opds[0]);
1748 break;
1749 case spirv_op_specconstantfalse:
1750 out("%%%u = %s type=%%%u\n", opds[1], op_name, opds[0]);
1751 break;
1752 case spirv_op_specconstant:
1753 op_specconstant_out();
1754 break;
1755 case spirv_op_specconstantcomposite:
1756 op_specconstantcomposite_out();
1757 break;
1758 case spirv_op_specconstantop:
1759 op_specconstantop_out();
1760 break;
1761 /* constants end ---------------------------------------------*/
1762 case spirv_op_variable:
1763 op_variable_out(false);
1764 break;
1765 case spirv_op_undef:
1766 out("%%%u = %s type=%%%u\n", opds[1], op_name, opds[0]);
1767 break;
1768 default:
1769 section_breaking_op_num = op_num;
1770 break;
1771 }
1772
1773 if (section_breaking_op_num != spirv_op_invalid)
1774 break;
1775
1776 r = read_whole_op();
1777 if (r != OK)
1778 break;
1779 }
1780 if (section_breaking_op_num == spirv_op_invalid) {
1781 out("// section end: non function declarations\n");
1782 exit(0);
1783 }
1784 out("// section end: non function declarations, breaking opcode=%s(%u)\n", ops_name[section_breaking_op_num], section_breaking_op_num);
1785 }
1786
1787 #define flag(val, str) \
1788 if ((w & val) != 0) { \
1789 if (!first) \
1790 strcat(b, "|"); \
1791 strcat(b, str); \
1792 first = false; \
1793 }
1794 static u8 *function_control_to_str(u32 w)
1795 {
1796 static u8 b[256];
1797 bool first;
1798
1799 if (w == 0)
1800 return "none";
1801
1802 first = true;
1803 b[0] = 0;
1804
1805 flag(0x00000001, "inline")
1806 flag(0x00000002, "dont_inline")
1807 flag(0x00000004, "pure")
1808 flag(0x00000008, "const")
1809 flag(0xfffffff0, "unknown_function_control_flag(s)")
1810 }
1811 #undef flag
1812
1813 static u8 *scope_to_str(u32 w)
1814 {
1815 switch (w) {
1816 case 0:
1817 return "cross_device";
1818 case 1:
1819 return "device";
1820 case 2:
1821 return "workgroup";
1822 case 3:
1823 return "subgroup";
1824 case 4:
1825 return "invocation";
1826 case 5:
1827 return "queue_family";
1828 default:
1829 return "unknown_scope_code";
1830 }
1831 }
1832
1833 #define flag_no_opd(val, str) \
1834 if ((opds[i] & val) != 0) { \
1835 if (!first) \
1836 out("|"); \
1837 out(str); \
1838 first = false; \
1839 }
1840 #define TMP_SZ 128
1841 /* return the idx of the next mem opd */
1842 static u16 mem_opd_out(u16 i)
1843 {
1844 u8 tmp[TMP_SZ];
1845 u16 additional_opds;
1846 bool first;
1847
1848 additional_opds = i + 1;
1849
1850 if (opds[i] == 0) {
1851 out("none");
1852 return i + 1;
1853 }
1854
1855 first = true;
1856
1857 /* order matters */
1858 flag_no_opd(0x00000001, "volatile")
1859
1860 if ((opds[i] & 0x00000002) != 0) {
1861 if (!first)
1862 out("|");
1863 out("aligned");
1864 first = false;
1865 snprintf(tmp, TMP_SZ, "(%u)", opds[additional_opds]);
1866 out(tmp);
1867 ++additional_opds;
1868 }
1869
1870 flag_no_opd(0x00000004, "non_temporal")
1871
1872 if ((opds[i] & 0x00000008) != 0) {
1873 if (!first)
1874 out("|");
1875 out("make_pointer_available");
1876 first = false;
1877 snprintf(tmp, TMP_SZ, "(%s)", scope_to_str(opds[additional_opds]));
1878 out(tmp);
1879 ++additional_opds;
1880 }
1881
1882 if ((opds[i] & 0x00000010) != 0) {
1883 if (!first)
1884 out("|");
1885 out("make_pointer_visible");
1886 first = false;
1887 snprintf(tmp, TMP_SZ, "(%s)", scope_to_str(opds[additional_opds]));
1888 out(tmp);
1889 ++additional_opds;
1890 }
1891
1892 flag_no_opd(0x00000010, "non_private_pointer")
1893
1894 flag_no_opd(0xffffffe0, "unknown_memory_operand_flag(s)-->consider the following instruction operands as corrupted")
1895
1896 return additional_opds;
1897 }
1898 #undef flag_no_opd
1899 #undef TMP_SZ
1900
1901 /* will output the mem opds till the end of the instruction */
1902 static void mem_opds_out(u16 i)
1903 {
1904 u16 mem_opd_idx;
1905
1906 mem_opd_idx = 0;
1907 out(" memory_operands[%u]=", mem_opd_idx);
1908 loop {
1909 i = mem_opd_out(i);
1910 ++mem_opd_idx;
1911
1912 if (i > opds_last)
1913 break;
1914 }
1915 }
1916
1917 static void op_load_out_depth(void)
1918 {
1919 out_depth("%%%u = %s type=%%%u pointer=%%%u", opds[1], op_name, opds[0], opds[2]);
1920 if (op_ws_n > 4)
1921 mem_opds_out(3);
1922 out("\n");
1923 }
1924
1925 static void op_vectorshuffle_out_depth(void)
1926 {
1927 u16 i;
1928
1929 out_depth("%%%u = %s type=%%%u vector_0=%%%u vector_1=%%%u", opds[1], op_name, opds[0], opds[2], opds[3]);
1930
1931 i = 4;
1932 loop {
1933 if (i > opds_last)
1934 break;
1935 out(" components[%u]=%u", i - 4, opds[i]);
1936 ++i;
1937 }
1938 out("\n");
1939 }
1940
1941 static void op_compositeextract_out_depth(void)
1942 {
1943 u16 i;
1944
1945 out_depth("%%%u = %s type=%%%u composite=%%%u", opds[1], op_name, opds[0], opds[2]);
1946
1947 i = 3;
1948 loop {
1949 if (i > opds_last)
1950 break;
1951 out(" indexes[%u]=%u", i - 3, opds[i]);
1952 ++i;
1953 }
1954 out("\n");
1955 }
1956
1957 static void op_compositeconstruct_out_depth(void)
1958 {
1959 u16 i;
1960
1961 out_depth("%%%u = %s type=%%%u", opds[1], op_name, opds[0]);
1962
1963 i = 2;
1964 loop {
1965 if (i > opds_last)
1966 break;
1967 out(" constituents[%u]=%%%u", i - 2, opds[i]);
1968 ++i;
1969 }
1970 out("\n");
1971 }
1972
1973 static void op_store_out_depth(void)
1974 {
1975 out_depth("%s pointer=%%%u object=%%%u", op_name, opds[0], opds[1]);
1976 if (op_ws_n > 3)
1977 mem_opds_out(2);
1978 out("\n");
1979 }
1980
1981 static void op_imagesampleimplicitlod_out_depth(void)
1982 {
1983 u16 i;
1984 out_depth("%%%u = %s type=%%%u sampled_image=%%%u coordinate=%%%u", opds[1], op_name, opds[0], opds[2], opds[3]);
1985
1986 /*
1987 * XXX: this instruction encoding seems serevely broken since it
1988 * seems to depend on previous instructions, or I do not
1989 * understand how to properly decode it yet
1990 */
1991 i = 4;
1992 loop {
1993 if (i > opds_last)
1994 break;
1995 out(" image_operands[%u]=0x%08x", i - 4, opds[i]);
1996 ++i;
1997 }
1998 out("\n");
1999 }
2000
2001 static void op_imagesampleexplicitlod_out_depth(void)
2002 {
2003 u16 i;
2004
2005 out_depth("%%%u = %s type=%%%u sampled_image=%%%u coordinate=%%%u", opds[1], op_name, opds[0], opds[2], opds[3]);
2006
2007 /*
2008 * XXX: this instruction encoding seems serevely broken since it
2009 * seems to depend on previous instructions, or I do not
2010 * understand how to properly decode it yet
2011 */
2012 i = 4;
2013 loop {
2014 if (i > opds_last)
2015 break;
2016 out(" image_operands[%u]=0x%08x", i - 4, opds[i]);
2017 ++i;
2018 }
2019 out("\n");
2020 }
2021
2022 static void op_functioncall_out_depth(void)
2023 {
2024 u16 i;
2025
2026 out_depth("%%%u = %s type=%%%u function=%%%u", opds[1], op_name, opds[0], opds[2]);
2027
2028 i = 3;
2029 loop {
2030 if (i > opds_last)
2031 break;
2032 out(" arguments[%u]=%%%u", i - 3, opds[i]);
2033 ++i;
2034 }
2035 out("\n");
2036 }
2037
2038 static void op_branchconditional_out_depth(void)
2039 {
2040 u16 i;
2041
2042 out_depth("%s condition=%%%u true_label=%%%u false_label %%%u", op_name, opds[0], opds[1], opds[2]);
2043
2044 i = 3;
2045 loop { /* 0 or 2 */
2046 if (i > opds_last)
2047 break;
2048 out(" weights[%u]=%u", i - 3, opds[i]);
2049 ++i;
2050 }
2051 out("\n");
2052 }
2053
2054 static u8 *selection_control_to_str(u32 w)
2055 {
2056 switch (w) {
2057 case 0:
2058 return "none";
2059 case 1:
2060 return "flatten";
2061 case 2:
2062 return "dont_flatten";
2063 default:
2064 return "unknown_selection_control_code";
2065 }
2066 }
2067
2068 /* TODO: add a extinst disassembler */
2069 static void op_extinst_out_depth(void)
2070 {
2071 u16 i;
2072
2073 out_depth("%%%u = %s type=%%%u set=%%%u instruction=%u", opds[1], op_name, opds[0], opds[2], opds[3]);
2074 i = 4;
2075 loop {
2076 if (i > opds_last)
2077 break;
2078 out(" operands[%u]=%%%u", i - 4, opds[i]);
2079 ++i;
2080 }
2081 out("\n");
2082 }
2083
2084 static void op_accesschain_out_depth(void)
2085 {
2086 u16 i;
2087
2088 out_depth("%%%u = %s type=%%%u base=%%%u", opds[1], op_name, opds[0], opds[2]);
2089
2090 i = 3;
2091 loop {
2092 if (i > opds_last)
2093 break;
2094 out(" indexes[%u]=%%%u", i - 3, opds[i]);
2095 ++i;
2096 }
2097 out("\n");
2098 }
2099
2100 static void op_compositeinsert_out_depth(void)
2101 {
2102 u16 i;
2103
2104 out_depth("%%%u = %s type=%%%u object=%%%u composite=%%%u", opds[1], op_name, opds[0], opds[2], opds[3]);
2105
2106 i = 4;
2107 loop {
2108 if (i > opds_last)
2109 break;
2110 out(" indexes[%u]=%u", i - 4, opds[i]);
2111 ++i;
2112 }
2113 out("\n");
2114 }
2115
2116 static void op_phi_out_depth(void)
2117 {
2118 u16 i;
2119
2120 out_depth("%%%u = %s type=%%%u", opds[1], op_name, opds[0]);
2121
2122 i = 2;
2123 loop {
2124 if (i > opds_last)
2125 break;
2126 out( "variables[%u]=%%%u parents[%u]=%%%u", (i - 2)/ 2, opds[i], opds[i + 1]);
2127 i += 2;
2128 }
2129 out("\n");
2130 }
2131
2132 /*
2133 * we do a bit of state tracking in order to detect function declarations
2134 * once function definitions did start, which is not allowed
2135 */
2136 static void layout_funcs_out(void)
2137 {
2138 u16 section_breaking_op_num;
2139 bool defs_section;
2140 bool func_has_blk;
2141
2142 defs_section = false;
2143
2144 section_breaking_op_num = spirv_op_invalid;
2145 out("\n// section start: function declarations then definitions\n");
2146 loop {
2147 u8 r;
2148
2149 switch (op_num) {
2150 case spirv_op_line:
2151 out_depth("%s %%%u line=%u column=%u\n", op_name, opds[0], opds[1], opds[2]);
2152 break;
2153 case spirv_op_function:
2154 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]);
2155 func_has_blk = false;
2156 ++depth; /* start of a blk */
2157 break;
2158 case spirv_op_functionparameter:
2159 out_depth("%%%u = %s type=%%%u\n", opds[1], op_name, opds[0]);
2160 break;
2161 case spirv_op_functionend:
2162 depth--; /* end of blk */
2163 out_depth("%s\n", op_name);
2164 if (!func_has_blk && defs_section) {
2165 out_depth("error: function declaration in function definition section\n");
2166 exit(1);
2167 }
2168 break;
2169 case spirv_op_label:
2170 func_has_blk = true;
2171 defs_section = true;
2172 out_depth("%%%u = %s\n", opds[0], op_name);
2173 ++depth; /* start of a blk */
2174 break;
2175 case spirv_op_load:
2176 op_load_out_depth();
2177 break;
2178 case spirv_op_vectorshuffle:
2179 op_vectorshuffle_out_depth();
2180 break;
2181 case spirv_op_compositeextract:
2182 op_compositeextract_out_depth();
2183 break;
2184 case spirv_op_compositeconstruct:
2185 op_compositeconstruct_out_depth();
2186 break;
2187 case spirv_op_store:
2188 op_store_out_depth();
2189 break;
2190 case spirv_op_return:
2191 out_depth("%s\n", op_name);
2192 --depth; /* end of blk */
2193 break;
2194 case spirv_op_imagesampleimplicitlod:
2195 op_imagesampleimplicitlod_out_depth();
2196 break;
2197 case spirv_op_fmul:
2198 out_depth("%%%u = %s type=%%%u operands[0]=%%%u operands[1]=%%%u\n", opds[1], op_name, opds[0], opds[2], opds[3]);
2199 break;
2200 case spirv_op_variable:
2201 op_variable_out(true);
2202 break;
2203 case spirv_op_functioncall:
2204 op_functioncall_out_depth();
2205 break;
2206 case spirv_op_fordlessthanequal:
2207 out_depth("%%%u = %s type=%%%u operands[0]=%%%u operands[1]=%%%u\n", opds[1], op_name, opds[0], opds[2], opds[3]);
2208 break;
2209 case spirv_op_selectionmerge:
2210 out_depth("%s merge_block=%%%u section_control=%s\n", op_name, opds[0], selection_control_to_str(opds[1]));
2211 break;
2212 case spirv_op_branchconditional:
2213 op_branchconditional_out_depth();
2214 depth--; /* end of blk */
2215 break;
2216 case spirv_op_fdiv:
2217 out_depth("%%%u = %s type=%%%u operands[0]=%%%u operands[1]=%%%u\n", opds[1], op_name, opds[0], opds[2], opds[3]);
2218 break;
2219 case spirv_op_branch:
2220 out_depth("%s target_label=%%%u\n", op_name, opds[0]);
2221 depth--; /* end of blk */
2222 break;
2223 case spirv_op_fadd:
2224 out_depth("%%%u = %s type=%%%u operands[0]=%%%u operands[1]=%%%u\n", opds[1], op_name, opds[0], opds[2], opds[3]);
2225 break;
2226 case spirv_op_extinst:
2227 op_extinst_out_depth();
2228 break;
2229 case spirv_op_returnvalue:
2230 out_depth("%s value=%%%u\n", op_name, opds[0]);
2231 --depth; /* end of blk */
2232 break;
2233 case spirv_op_accesschain:
2234 op_accesschain_out_depth();
2235 break;
2236 case spirv_op_compositeinsert:
2237 op_compositeinsert_out_depth();
2238 break;
2239 case spirv_op_sampledimage:
2240 out_depth("%%%u = %s type=%%%u image=%%%u sampler=%%%u\n", opds[1], op_name, opds[0], opds[2], opds[3]);
2241 break;
2242 case spirv_op_negate:
2243 out_depth("%%%u = %s type=%%%u operand=%%%u\n", opds[1], op_name, opds[0], opds[2]);
2244 break;
2245 case spirv_op_vectortimesscalar:
2246 out_depth("%%%u = %s type=%%%u vector=%%%u scalar=%%%u\n", opds[1], op_name, opds[0], opds[2], opds[3]);
2247 break;
2248 case spirv_op_fsub:
2249 out_depth("%%%u = %s type=%%%u operands[0]=%%%u operands[1]=%%%u\n", opds[1], op_name, opds[0], opds[2], opds[3]);
2250 break;
2251 case spirv_op_fordnotequal:
2252 out_depth("%%%u = %s type=%%%u operands[0]=%%%u operands[1]=%%%u\n", opds[1], op_name, opds[0], opds[2], opds[3]);
2253 break;
2254 case spirv_op_fordgreaterthanequal:
2255 out_depth("%%%u = %s type=%%%u operands[0]=%%%u operands[1]=%%%u\n", opds[1], op_name, opds[0], opds[2], opds[3]);
2256 break;
2257 case spirv_op_phi:
2258 op_phi_out_depth();
2259 break;
2260 case spirv_op_fordgreaterthan:
2261 out_depth("%%%u = %s type=%%%u operands[0]=%%%u operands[1]=%%%u\n", opds[1], op_name, opds[0], opds[2], opds[3]);
2262 break;
2263 case spirv_op_logicalor:
2264 out_depth("%%%u = %s type=%%%u operands[0]=%%%u operands[1]=%%%u\n", opds[1], op_name, opds[0], opds[2], opds[3]);
2265 break;
2266 case spirv_op_fordequal:
2267 out_depth("%%%u = %s type=%%%u operands[0]=%%%u operands[1]=%%%u\n", opds[1], op_name, opds[0], opds[2], opds[3]);
2268 break;
2269 case spirv_op_select:
2270 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]);
2271 break;
2272 case spirv_op_imagesampleexplicitlod:
2273 op_imagesampleexplicitlod_out_depth();
2274 break;
2275 case spirv_op_fordlessthan:
2276 out_depth("%%%u = %s type=%%%u operands[0]=%%%u operands[1]=%%%u\n", opds[1], op_name, opds[0], opds[2], opds[3]);
2277 break;
2278 case spirv_op_vectortimesmatrix:
2279 out_depth("%%%u = %s type=%%%u vector=%%%u matrix=%%%u\n", opds[1], op_name, opds[0], opds[2], opds[3]);
2280 break;
2281 case spirv_op_dot:
2282 out_depth("%%%u = %s type=%%%u vectors[0]=%%%u vectors[1]=%%%u\n", opds[1], op_name, opds[0], opds[2], opds[3]);
2283 break;
2284 case spirv_op_fmod:
2285 out_depth("%%%u = %s type=%%%u operands[0]=%%%u operands[1]=%%%u\n", opds[1], op_name, opds[0], opds[2], opds[3]);
2286 break;
2287 case spirv_op_logicaland:
2288 out_depth("%%%u = %s type=%%%u operands[0]=%%%u operands[1]=%%%u\n", opds[1], op_name, opds[0], opds[2], opds[3]);
2289 break;
2290 /* TODO: MORE! */
2291 default:
2292 section_breaking_op_num = op_num;
2293 break;
2294 }
2295
2296 if (section_breaking_op_num != spirv_op_invalid)
2297 break;
2298
2299 r = read_whole_op();
2300 if (r != OK)
2301 break;
2302 }
2303 if (section_breaking_op_num == spirv_op_invalid) {
2304 out("// section end: function declarations then definitions\n");
2305 exit(0);
2306 }
2307 out("// section end: function declarations and definitions, breaking opcode=%s(%u)\n", ops_name[section_breaking_op_num], section_breaking_op_num);
2308 }
2309
2310 static void init_misc(void)
2311 {
2312 u16 i;
2313
2314 i = 0;
2315 loop {
2316 if (i == 0xffff)
2317 break;
2318 if (ops_name[i] == 0)
2319 ops_name[i] = "unkwown_opcode";
2320 ++i;
2321 }
2322
2323 depth = 0;
2324 }
2325
2326 int main(void)
2327 {
2328 u8 r;
2329
2330 init_misc();
2331
2332 clearerr(stdin);
2333 hdr_out();
2334
2335 r = read_whole_op();
2336 if (r != OK)
2337 exit(0);
2338
2339 /* logical layout of a spirv module ----------------------------------*/
2340 layout_caps_out();
2341 layout_extns_out();
2342 layout_extinstimports_out();
2343 layout_memorymodel_out();
2344 layout_entrypoints_out();
2345 layout_executionmodes_out();
2346 layout_debug_out();
2347 layout_annotation_out();
2348 /* from here opline is allowed */
2349 layout_nonfuncdecls_out();
2350 layout_funcs_out();
2351 /* -------------------------------------------------------------------*/
2352 return 0;
2353 }
Hints:
Before first commit, do not forget to setup your git environment:
git config --global user.name "your_name_here"
git config --global user.email "your@email_here"

Clone this repository using HTTP(S):
git clone https://rocketgit.com/user/sylware/vulkan-misc

Clone this repository using ssh (do not forget to upload a key first):
git clone ssh://rocketgit@ssh.rocketgit.com/user/sylware/vulkan-misc

Clone this repository using git:
git clone git://git.rocketgit.com/user/sylware/vulkan-misc

You are allowed to anonymously push to this repository.
This means that your pushed commits will automatically be transformed into a merge request:
... clone the repository ...
... make some changes and some commits ...
git push origin main