1 /*
2 * Copyright © 2014 Intel Corporation
3 *
4 * Permission is hereby granted, free of charge, to any person obtaining a
5 * copy of this software and associated documentation files (the "Software"),
6 * to deal in the Software without restriction, including without limitation
7 * the rights to use, copy, modify, merge, publish, distribute, sublicense,
8 * and/or sell copies of the Software, and to permit persons to whom the
9 * Software is furnished to do so, subject to the following conditions:
10 *
11 * The above copyright notice and this permission notice (including the next
12 * paragraph) shall be included in all copies or substantial portions of the
13 * Software.
14 *
15 * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
16 * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
17 * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL
18 * THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
19 * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING
20 * FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS
21 * IN THE SOFTWARE.
22 *
23 * Authors:
24 * Connor Abbott (cwabbott0@gmail.com)
25 *
26 */
27
28 #include <inttypes.h> /* for PRIx64 macro */
29 #include <math.h>
30 #include <stdio.h>
31 #include <stdlib.h>
32 #include "compiler/shader_enums.h"
33 #include "util/half_float.h"
34 #include "util/memstream.h"
35 #include "util/mesa-blake3.h"
36 #include "vulkan/vulkan_core.h"
37 #include "nir.h"
38 #include "nir_builder.h"
39
40 static void
print_indentation(unsigned levels,FILE * fp)41 print_indentation(unsigned levels, FILE *fp)
42 {
43 for (unsigned i = 0; i < levels; i++)
44 fprintf(fp, " ");
45 }
46
47 typedef struct {
48 FILE *fp;
49 nir_shader *shader;
50
51 const char *def_prefix;
52
53 /** map from nir_variable -> printable name */
54 struct hash_table *ht;
55
56 /** set of names used so far for nir_variables */
57 struct set *syms;
58
59 /* an index used to make new non-conflicting names */
60 unsigned index;
61
62 /* Used with nir_gather_types() to identify best representation
63 * to print terse inline constant values together with SSA sources.
64 * Updated per nir_function_impl being printed.
65 */
66 BITSET_WORD *float_types;
67 BITSET_WORD *int_types;
68
69 /**
70 * Optional table of annotations mapping nir object
71 * (such as instr or var) to message to print.
72 */
73 struct hash_table *annotations;
74
75 /* Maximum length for SSA or Reg index in the current impl */
76 unsigned max_dest_index;
77
78 /* Padding for instructions without destination to make
79 * them align with the `=` for instructions with destination.
80 */
81 unsigned padding_for_no_dest;
82
83 nir_debug_info_instr **debug_info;
84 } print_state;
85
86 static void
print_annotation(print_state * state,void * obj)87 print_annotation(print_state *state, void *obj)
88 {
89 FILE *fp = state->fp;
90
91 if (!state->annotations)
92 return;
93
94 struct hash_entry *entry = _mesa_hash_table_search(state->annotations, obj);
95 if (!entry)
96 return;
97
98 const char *note = entry->data;
99 _mesa_hash_table_remove(state->annotations, entry);
100
101 fprintf(fp, "%s\n\n", note);
102 }
103
104 /* For 1 element, the size is intentionally omitted. */
105 static const char *sizes[] = { "x??", " ", "x2 ", "x3 ", "x4 ",
106 "x5 ", "x??", "x??", "x8 ",
107 "x??", "x??", "x??", "x??",
108 "x??", "x??", "x??", "x16" };
109
110 static const char *
divergence_status(print_state * state,bool divergent)111 divergence_status(print_state *state, bool divergent)
112 {
113 if (state->shader->info.divergence_analysis_run)
114 return divergent ? "div " : "con ";
115
116 return "";
117 }
118
119 static unsigned
count_digits(unsigned n)120 count_digits(unsigned n)
121 {
122 return n ? (unsigned)floor(log10(n)) + 1u : 1u;
123 }
124
125 static void
print_def(nir_def * def,print_state * state)126 print_def(nir_def *def, print_state *state)
127 {
128 FILE *fp = state->fp;
129
130 const unsigned ssa_padding = state->max_dest_index ? count_digits(state->max_dest_index) - count_digits(def->index) : 0;
131
132 const unsigned padding = (def->bit_size <= 8) + 1 + ssa_padding;
133
134 fprintf(fp, "%s%u%s%*s%s%u",
135 divergence_status(state, def->divergent),
136 def->bit_size, sizes[def->num_components],
137 padding, "", state->def_prefix, def->index);
138 }
139
140 static unsigned
calculate_padding_for_no_dest(print_state * state)141 calculate_padding_for_no_dest(print_state *state)
142 {
143 const unsigned div = state->shader->info.divergence_analysis_run ? 4 : 0;
144 const unsigned ssa_size = 5;
145 const unsigned percent = 1;
146 const unsigned ssa_index = count_digits(state->max_dest_index);
147 const unsigned equals = 1;
148 return ssa_size + 1 + div + percent + ssa_index + 1 + equals + 1;
149 }
150
151 static void
print_no_dest_padding(print_state * state)152 print_no_dest_padding(print_state *state)
153 {
154 FILE *fp = state->fp;
155
156 if (state->padding_for_no_dest)
157 fprintf(fp, "%*s", state->padding_for_no_dest, "");
158 }
159
160 static void
print_hex_padded_const_value(const nir_const_value * value,unsigned bit_size,FILE * fp)161 print_hex_padded_const_value(const nir_const_value *value, unsigned bit_size, FILE *fp)
162 {
163 switch (bit_size) {
164 case 64:
165 fprintf(fp, "0x%016" PRIx64, value->u64);
166 break;
167 case 32:
168 fprintf(fp, "0x%08x", value->u32);
169 break;
170 case 16:
171 fprintf(fp, "0x%04x", value->u16);
172 break;
173 case 8:
174 fprintf(fp, "0x%02x", value->u8);
175 break;
176 default:
177 unreachable("unhandled bit size");
178 }
179 }
180
181 static void
print_hex_terse_const_value(const nir_const_value * value,unsigned bit_size,FILE * fp)182 print_hex_terse_const_value(const nir_const_value *value, unsigned bit_size, FILE *fp)
183 {
184 switch (bit_size) {
185 case 64:
186 fprintf(fp, "0x%" PRIx64, value->u64);
187 break;
188 case 32:
189 fprintf(fp, "0x%x", value->u32);
190 break;
191 case 16:
192 fprintf(fp, "0x%x", value->u16);
193 break;
194 case 8:
195 fprintf(fp, "0x%x", value->u8);
196 break;
197 default:
198 unreachable("unhandled bit size");
199 }
200 }
201
202 static void
print_float_const_value(const nir_const_value * value,unsigned bit_size,FILE * fp)203 print_float_const_value(const nir_const_value *value, unsigned bit_size, FILE *fp)
204 {
205 switch (bit_size) {
206 case 64:
207 fprintf(fp, "%f", value->f64);
208 break;
209 case 32:
210 fprintf(fp, "%f", value->f32);
211 break;
212 case 16:
213 fprintf(fp, "%f", _mesa_half_to_float(value->u16));
214 break;
215 default:
216 unreachable("unhandled bit size");
217 }
218 }
219
220 static void
print_int_const_value(const nir_const_value * value,unsigned bit_size,FILE * fp)221 print_int_const_value(const nir_const_value *value, unsigned bit_size, FILE *fp)
222 {
223 switch (bit_size) {
224 case 64:
225 fprintf(fp, "%+" PRIi64, value->i64);
226 break;
227 case 32:
228 fprintf(fp, "%+d", value->i32);
229 break;
230 case 16:
231 fprintf(fp, "%+d", value->i16);
232 break;
233 case 8:
234 fprintf(fp, "%+d", value->i8);
235 break;
236 default:
237 unreachable("unhandled bit size");
238 }
239 }
240
241 static void
print_uint_const_value(const nir_const_value * value,unsigned bit_size,FILE * fp)242 print_uint_const_value(const nir_const_value *value, unsigned bit_size, FILE *fp)
243 {
244 switch (bit_size) {
245 case 64:
246 fprintf(fp, "%" PRIu64, value->u64);
247 break;
248 case 32:
249 fprintf(fp, "%u", value->u32);
250 break;
251 case 16:
252 fprintf(fp, "%u", value->u16);
253 break;
254 case 8:
255 fprintf(fp, "%u", value->u8);
256 break;
257 default:
258 unreachable("unhandled bit size");
259 }
260 }
261
262 static void
print_const_from_load(nir_load_const_instr * instr,print_state * state,nir_alu_type type)263 print_const_from_load(nir_load_const_instr *instr, print_state *state, nir_alu_type type)
264 {
265 FILE *fp = state->fp;
266
267 const unsigned bit_size = instr->def.bit_size;
268 const unsigned num_components = instr->def.num_components;
269
270 type = nir_alu_type_get_base_type(type);
271
272 /* There's only one way to print booleans. */
273 if (bit_size == 1 || type == nir_type_bool) {
274 fprintf(fp, "(");
275 for (unsigned i = 0; i < num_components; i++) {
276 if (i != 0)
277 fprintf(fp, ", ");
278 fprintf(fp, "%s", instr->value[i].b ? "true" : "false");
279 }
280 fprintf(fp, ")");
281 return;
282 }
283
284 fprintf(fp, "(");
285
286 if (type != nir_type_invalid) {
287 for (unsigned i = 0; i < num_components; i++) {
288 const nir_const_value *v = &instr->value[i];
289 if (i != 0)
290 fprintf(fp, ", ");
291 switch (type) {
292 case nir_type_float:
293 print_float_const_value(v, bit_size, fp);
294 break;
295 case nir_type_int:
296 case nir_type_uint:
297 print_hex_terse_const_value(v, bit_size, fp);
298 break;
299
300 default:
301 unreachable("invalid nir alu base type");
302 }
303 }
304 } else {
305 #define PRINT_VALUES(F) \
306 do { \
307 for (unsigned i = 0; i < num_components; i++) { \
308 if (i != 0) \
309 fprintf(fp, ", "); \
310 F(&instr->value[i], bit_size, fp); \
311 } \
312 } while (0)
313
314 #define SEPARATOR() \
315 if (num_components > 1) \
316 fprintf(fp, ") = ("); \
317 else \
318 fprintf(fp, " = ")
319
320 bool needs_float = bit_size > 8;
321 bool needs_signed = false;
322 bool needs_decimal = false;
323 for (unsigned i = 0; i < num_components; i++) {
324 const nir_const_value *v = &instr->value[i];
325 switch (bit_size) {
326 case 64:
327 needs_signed |= v->i64 < 0;
328 needs_decimal |= v->u64 >= 10;
329 break;
330 case 32:
331 needs_signed |= v->i32 < 0;
332 needs_decimal |= v->u32 >= 10;
333 break;
334 case 16:
335 needs_signed |= v->i16 < 0;
336 needs_decimal |= v->u16 >= 10;
337 break;
338 case 8:
339 needs_signed |= v->i8 < 0;
340 needs_decimal |= v->u8 >= 10;
341 break;
342 default:
343 unreachable("invalid bit size");
344 }
345 }
346
347 if (state->int_types) {
348 const unsigned index = instr->def.index;
349 const bool inferred_int = BITSET_TEST(state->int_types, index);
350 const bool inferred_float = BITSET_TEST(state->float_types, index);
351
352 if (inferred_int && !inferred_float) {
353 needs_float = false;
354 } else if (inferred_float && !inferred_int) {
355 needs_signed = false;
356 needs_decimal = false;
357 }
358 }
359
360 PRINT_VALUES(print_hex_padded_const_value);
361
362 if (needs_float) {
363 SEPARATOR();
364 PRINT_VALUES(print_float_const_value);
365 }
366
367 if (needs_signed) {
368 SEPARATOR();
369 PRINT_VALUES(print_int_const_value);
370 }
371
372 if (needs_decimal) {
373 SEPARATOR();
374 PRINT_VALUES(print_uint_const_value);
375 }
376 }
377
378 fprintf(fp, ")");
379 }
380
381 static void
print_load_const_instr(nir_load_const_instr * instr,print_state * state)382 print_load_const_instr(nir_load_const_instr *instr, print_state *state)
383 {
384 FILE *fp = state->fp;
385
386 print_def(&instr->def, state);
387
388 fprintf(fp, " = load_const ");
389
390 /* In the definition, print all interpretations of the value. */
391 print_const_from_load(instr, state, nir_type_invalid);
392 }
393
394 static void
print_src(const nir_src * src,print_state * state,nir_alu_type src_type)395 print_src(const nir_src *src, print_state *state, nir_alu_type src_type)
396 {
397 FILE *fp = state->fp;
398 fprintf(fp, "%s%u", state->def_prefix, src->ssa->index);
399 nir_instr *instr = src->ssa->parent_instr;
400
401 if (instr->type == nir_instr_type_load_const && !NIR_DEBUG(PRINT_NO_INLINE_CONSTS)) {
402 nir_load_const_instr *load_const = nir_instr_as_load_const(instr);
403 fprintf(fp, " ");
404
405 nir_alu_type type = nir_alu_type_get_base_type(src_type);
406
407 if (type == nir_type_invalid && state->int_types) {
408 const unsigned index = load_const->def.index;
409 const bool inferred_int = BITSET_TEST(state->int_types, index);
410 const bool inferred_float = BITSET_TEST(state->float_types, index);
411
412 if (inferred_float && !inferred_int)
413 type = nir_type_float;
414 }
415
416 if (type == nir_type_invalid)
417 type = nir_type_uint;
418
419 /* For a constant in a source, always pick one interpretation. */
420 assert(type != nir_type_invalid);
421 print_const_from_load(load_const, state, type);
422 }
423 }
424
425 static const char *
comp_mask_string(unsigned num_components)426 comp_mask_string(unsigned num_components)
427 {
428 return (num_components > 4) ? "abcdefghijklmnop" : "xyzw";
429 }
430
431 static void
print_alu_src(nir_alu_instr * instr,unsigned src,print_state * state)432 print_alu_src(nir_alu_instr *instr, unsigned src, print_state *state)
433 {
434 FILE *fp = state->fp;
435
436 const nir_op_info *info = &nir_op_infos[instr->op];
437 print_src(&instr->src[src].src, state, info->input_types[src]);
438
439 bool print_swizzle = false;
440 nir_component_mask_t used_channels = 0;
441
442 for (unsigned i = 0; i < NIR_MAX_VEC_COMPONENTS; i++) {
443 if (!nir_alu_instr_channel_used(instr, src, i))
444 continue;
445
446 used_channels++;
447
448 if (instr->src[src].swizzle[i] != i) {
449 print_swizzle = true;
450 break;
451 }
452 }
453
454 unsigned live_channels = nir_src_num_components(instr->src[src].src);
455
456 if (print_swizzle || used_channels != live_channels) {
457 fprintf(fp, ".");
458 for (unsigned i = 0; i < NIR_MAX_VEC_COMPONENTS; i++) {
459 if (!nir_alu_instr_channel_used(instr, src, i))
460 continue;
461
462 fprintf(fp, "%c", comp_mask_string(live_channels)[instr->src[src].swizzle[i]]);
463 }
464 }
465 }
466
467 static void
print_alu_instr(nir_alu_instr * instr,print_state * state)468 print_alu_instr(nir_alu_instr *instr, print_state *state)
469 {
470 FILE *fp = state->fp;
471
472 print_def(&instr->def, state);
473
474 fprintf(fp, " = %s", nir_op_infos[instr->op].name);
475 if (instr->exact)
476 fprintf(fp, "!");
477 if (instr->no_signed_wrap)
478 fprintf(fp, ".nsw");
479 if (instr->no_unsigned_wrap)
480 fprintf(fp, ".nuw");
481 fprintf(fp, " ");
482
483 for (unsigned i = 0; i < nir_op_infos[instr->op].num_inputs; i++) {
484 if (i != 0)
485 fprintf(fp, ", ");
486
487 print_alu_src(instr, i, state);
488 }
489 }
490
491 static const char *
get_var_name(nir_variable * var,print_state * state)492 get_var_name(nir_variable *var, print_state *state)
493 {
494 if (state->ht == NULL)
495 return var->name ? var->name : "unnamed";
496
497 assert(state->syms);
498
499 struct hash_entry *entry = _mesa_hash_table_search(state->ht, var);
500 if (entry)
501 return entry->data;
502
503 char *name;
504 if (var->name == NULL) {
505 name = ralloc_asprintf(state->syms, "#%u", state->index++);
506 } else {
507 struct set_entry *set_entry = _mesa_set_search(state->syms, var->name);
508 if (set_entry != NULL) {
509 /* we have a collision with another name, append an # + a unique
510 * index */
511 name = ralloc_asprintf(state->syms, "%s#%u", var->name,
512 state->index++);
513 } else {
514 /* Mark this one as seen */
515 _mesa_set_add(state->syms, var->name);
516 name = var->name;
517 }
518 }
519
520 _mesa_hash_table_insert(state->ht, var, name);
521
522 return name;
523 }
524
525 static const char *
get_constant_sampler_addressing_mode(enum cl_sampler_addressing_mode mode)526 get_constant_sampler_addressing_mode(enum cl_sampler_addressing_mode mode)
527 {
528 switch (mode) {
529 case SAMPLER_ADDRESSING_MODE_NONE:
530 return "none";
531 case SAMPLER_ADDRESSING_MODE_CLAMP_TO_EDGE:
532 return "clamp_to_edge";
533 case SAMPLER_ADDRESSING_MODE_CLAMP:
534 return "clamp";
535 case SAMPLER_ADDRESSING_MODE_REPEAT:
536 return "repeat";
537 case SAMPLER_ADDRESSING_MODE_REPEAT_MIRRORED:
538 return "repeat_mirrored";
539 default:
540 unreachable("Invalid addressing mode");
541 }
542 }
543
544 static const char *
get_constant_sampler_filter_mode(enum cl_sampler_filter_mode mode)545 get_constant_sampler_filter_mode(enum cl_sampler_filter_mode mode)
546 {
547 switch (mode) {
548 case SAMPLER_FILTER_MODE_NEAREST:
549 return "nearest";
550 case SAMPLER_FILTER_MODE_LINEAR:
551 return "linear";
552 default:
553 unreachable("Invalid filter mode");
554 }
555 }
556
557 static void
print_constant(nir_constant * c,const struct glsl_type * type,print_state * state)558 print_constant(nir_constant *c, const struct glsl_type *type, print_state *state)
559 {
560 FILE *fp = state->fp;
561 const unsigned rows = glsl_get_vector_elements(type);
562 const unsigned cols = glsl_get_matrix_columns(type);
563 unsigned i;
564
565 switch (glsl_get_base_type(type)) {
566 case GLSL_TYPE_BOOL:
567 /* Only float base types can be matrices. */
568 assert(cols == 1);
569
570 for (i = 0; i < rows; i++) {
571 if (i > 0)
572 fprintf(fp, ", ");
573 fprintf(fp, "%s", c->values[i].b ? "true" : "false");
574 }
575 break;
576
577 case GLSL_TYPE_UINT8:
578 case GLSL_TYPE_INT8:
579 /* Only float base types can be matrices. */
580 assert(cols == 1);
581
582 for (i = 0; i < rows; i++) {
583 if (i > 0)
584 fprintf(fp, ", ");
585 fprintf(fp, "0x%02x", c->values[i].u8);
586 }
587 break;
588
589 case GLSL_TYPE_UINT16:
590 case GLSL_TYPE_INT16:
591 /* Only float base types can be matrices. */
592 assert(cols == 1);
593
594 for (i = 0; i < rows; i++) {
595 if (i > 0)
596 fprintf(fp, ", ");
597 fprintf(fp, "0x%04x", c->values[i].u16);
598 }
599 break;
600
601 case GLSL_TYPE_UINT:
602 case GLSL_TYPE_INT:
603 /* Only float base types can be matrices. */
604 assert(cols == 1);
605
606 for (i = 0; i < rows; i++) {
607 if (i > 0)
608 fprintf(fp, ", ");
609 fprintf(fp, "0x%08x", c->values[i].u32);
610 }
611 break;
612
613 case GLSL_TYPE_FLOAT16:
614 case GLSL_TYPE_FLOAT:
615 case GLSL_TYPE_DOUBLE:
616 if (cols > 1) {
617 for (i = 0; i < cols; i++) {
618 if (i > 0)
619 fprintf(fp, ", ");
620 print_constant(c->elements[i], glsl_get_column_type(type), state);
621 }
622 } else {
623 switch (glsl_get_base_type(type)) {
624 case GLSL_TYPE_FLOAT16:
625 for (i = 0; i < rows; i++) {
626 if (i > 0)
627 fprintf(fp, ", ");
628 fprintf(fp, "%f", _mesa_half_to_float(c->values[i].u16));
629 }
630 break;
631
632 case GLSL_TYPE_FLOAT:
633 for (i = 0; i < rows; i++) {
634 if (i > 0)
635 fprintf(fp, ", ");
636 fprintf(fp, "%f", c->values[i].f32);
637 }
638 break;
639
640 case GLSL_TYPE_DOUBLE:
641 for (i = 0; i < rows; i++) {
642 if (i > 0)
643 fprintf(fp, ", ");
644 fprintf(fp, "%f", c->values[i].f64);
645 }
646 break;
647
648 default:
649 unreachable("Cannot get here from the first level switch");
650 }
651 }
652 break;
653
654 case GLSL_TYPE_UINT64:
655 case GLSL_TYPE_INT64:
656 /* Only float base types can be matrices. */
657 assert(cols == 1);
658
659 for (i = 0; i < cols; i++) {
660 if (i > 0)
661 fprintf(fp, ", ");
662 fprintf(fp, "0x%08" PRIx64, c->values[i].u64);
663 }
664 break;
665
666 case GLSL_TYPE_STRUCT:
667 case GLSL_TYPE_INTERFACE:
668 for (i = 0; i < c->num_elements; i++) {
669 if (i > 0)
670 fprintf(fp, ", ");
671 fprintf(fp, "{ ");
672 print_constant(c->elements[i], glsl_get_struct_field(type, i), state);
673 fprintf(fp, " }");
674 }
675 break;
676
677 case GLSL_TYPE_ARRAY:
678 for (i = 0; i < c->num_elements; i++) {
679 if (i > 0)
680 fprintf(fp, ", ");
681 fprintf(fp, "{ ");
682 print_constant(c->elements[i], glsl_get_array_element(type), state);
683 fprintf(fp, " }");
684 }
685 break;
686
687 default:
688 unreachable("not reached");
689 }
690 }
691
692 static const char *
get_variable_mode_str(nir_variable_mode mode,bool want_local_global_mode)693 get_variable_mode_str(nir_variable_mode mode, bool want_local_global_mode)
694 {
695 switch (mode) {
696 case nir_var_shader_in:
697 return "shader_in";
698 case nir_var_shader_out:
699 return "shader_out";
700 case nir_var_uniform:
701 return "uniform";
702 case nir_var_mem_ubo:
703 return "ubo";
704 case nir_var_system_value:
705 return "system";
706 case nir_var_mem_ssbo:
707 return "ssbo";
708 case nir_var_mem_shared:
709 return "shared";
710 case nir_var_mem_global:
711 return "global";
712 case nir_var_mem_push_const:
713 return "push_const";
714 case nir_var_mem_constant:
715 return "constant";
716 case nir_var_image:
717 return "image";
718 case nir_var_shader_temp:
719 return want_local_global_mode ? "shader_temp" : "";
720 case nir_var_function_temp:
721 return want_local_global_mode ? "function_temp" : "";
722 case nir_var_shader_call_data:
723 return "shader_call_data";
724 case nir_var_ray_hit_attrib:
725 return "ray_hit_attrib";
726 case nir_var_mem_task_payload:
727 return "task_payload";
728 case nir_var_mem_node_payload:
729 return "node_payload";
730 case nir_var_mem_node_payload_in:
731 return "node_payload_in";
732 default:
733 if (mode && (mode & nir_var_mem_generic) == mode)
734 return "generic";
735 return "";
736 }
737 }
738
739 static const char *
get_location_str(unsigned location,gl_shader_stage stage,nir_variable_mode mode,char * buf)740 get_location_str(unsigned location, gl_shader_stage stage,
741 nir_variable_mode mode, char *buf)
742 {
743 switch (stage) {
744 case MESA_SHADER_VERTEX:
745 if (mode == nir_var_shader_in)
746 return gl_vert_attrib_name(location);
747 else if (mode == nir_var_shader_out)
748 return gl_varying_slot_name_for_stage(location, stage);
749
750 break;
751 case MESA_SHADER_TESS_CTRL:
752 case MESA_SHADER_TESS_EVAL:
753 case MESA_SHADER_TASK:
754 case MESA_SHADER_MESH:
755 case MESA_SHADER_GEOMETRY:
756 if (mode == nir_var_shader_in || mode == nir_var_shader_out)
757 return gl_varying_slot_name_for_stage(location, stage);
758
759 break;
760 case MESA_SHADER_FRAGMENT:
761 if (mode == nir_var_shader_in)
762 return gl_varying_slot_name_for_stage(location, stage);
763 else if (mode == nir_var_shader_out)
764 return gl_frag_result_name(location);
765
766 break;
767 case MESA_SHADER_COMPUTE:
768 case MESA_SHADER_KERNEL:
769 default:
770 /* TODO */
771 break;
772 }
773
774 if (mode == nir_var_system_value)
775 return gl_system_value_name(location);
776
777 if (location == ~0) {
778 return "~0";
779 } else {
780 snprintf(buf, 4, "%u", location);
781 return buf;
782 }
783 }
784
785 static void
print_access(enum gl_access_qualifier access,print_state * state,const char * separator)786 print_access(enum gl_access_qualifier access, print_state *state, const char *separator)
787 {
788 if (!access) {
789 fputs("none", state->fp);
790 return;
791 }
792
793 static const struct {
794 enum gl_access_qualifier bit;
795 const char *name;
796 } modes[] = {
797 { ACCESS_COHERENT, "coherent" },
798 { ACCESS_RESTRICT, "restrict" },
799 { ACCESS_VOLATILE, "volatile" },
800 { ACCESS_NON_WRITEABLE, "readonly" },
801 { ACCESS_NON_READABLE, "writeonly" },
802 { ACCESS_NON_UNIFORM, "non-uniform" },
803 { ACCESS_CAN_REORDER, "reorderable" },
804 { ACCESS_NON_TEMPORAL, "non-temporal" },
805 { ACCESS_INCLUDE_HELPERS, "include-helpers" },
806 { ACCESS_IS_SWIZZLED_AMD, "is-swizzled-amd" },
807 { ACCESS_USES_FORMAT_AMD, "uses-format-amd" },
808 { ACCESS_FMASK_LOWERED_AMD, "fmask-lowered-amd" },
809 { ACCESS_CAN_SPECULATE, "speculatable" },
810 { ACCESS_CP_GE_COHERENT_AMD, "cp-ge-coherent-amd" },
811 { ACCESS_IN_BOUNDS_AGX, "in-bounds-agx" },
812 { ACCESS_KEEP_SCALAR, "keep-scalar" },
813 { ACCESS_SMEM_AMD, "smem-amd" },
814 };
815
816 bool first = true;
817 for (unsigned i = 0; i < ARRAY_SIZE(modes); ++i) {
818 if (access & modes[i].bit) {
819 fprintf(state->fp, "%s%s", first ? "" : separator, modes[i].name);
820 first = false;
821 }
822 }
823 }
824
825 static void
print_var_decl(nir_variable * var,print_state * state)826 print_var_decl(nir_variable *var, print_state *state)
827 {
828 FILE *fp = state->fp;
829
830 fprintf(fp, "decl_var ");
831
832 const char *const bindless = (var->data.bindless) ? "bindless " : "";
833 const char *const cent = (var->data.centroid) ? "centroid " : "";
834 const char *const samp = (var->data.sample) ? "sample " : "";
835 const char *const patch = (var->data.patch) ? "patch " : "";
836 const char *const inv = (var->data.invariant) ? "invariant " : "";
837 const char *const per_view = (var->data.per_view) ? "per_view " : "";
838 const char *const per_primitive = (var->data.per_primitive) ? "per_primitive " : "";
839 const char *const ray_query = (var->data.ray_query) ? "ray_query " : "";
840 const char *const fb_fetch = var->data.fb_fetch_output ? "fb_fetch_output " : "";
841 fprintf(fp, "%s%s%s%s%s%s%s%s%s%s %s ",
842 bindless, cent, samp, patch, inv, per_view, per_primitive,
843 ray_query, fb_fetch,
844 get_variable_mode_str(var->data.mode, false),
845 glsl_interp_mode_name(var->data.interpolation));
846
847 print_access(var->data.access, state, " ");
848 fprintf(fp, " ");
849
850 if (glsl_get_base_type(glsl_without_array(var->type)) == GLSL_TYPE_IMAGE) {
851 fprintf(fp, "%s ", util_format_short_name(var->data.image.format));
852 }
853
854 if (var->data.precision) {
855 const char *precisions[] = {
856 "",
857 "highp",
858 "mediump",
859 "lowp",
860 };
861 fprintf(fp, "%s ", precisions[var->data.precision]);
862 }
863
864 fprintf(fp, "%s %s", glsl_get_type_name(var->type),
865 get_var_name(var, state));
866
867 if (var->data.mode & (nir_var_shader_in |
868 nir_var_shader_out |
869 nir_var_uniform |
870 nir_var_system_value |
871 nir_var_mem_ubo |
872 nir_var_mem_ssbo |
873 nir_var_image)) {
874 char buf[4];
875 const char *loc = get_location_str(var->data.location,
876 state->shader->info.stage,
877 var->data.mode, buf);
878
879 /* For shader I/O vars that have been split to components or packed,
880 * print the fractional location within the input/output.
881 */
882 unsigned int num_components =
883 glsl_get_components(glsl_without_array(var->type));
884 const char *components = "";
885 char components_local[18] = { '.' /* the rest is 0-filled */ };
886 switch (var->data.mode) {
887 case nir_var_shader_in:
888 case nir_var_shader_out:
889 if (num_components < 16 && num_components != 0) {
890 const char *xyzw = comp_mask_string(num_components);
891 for (int i = 0; i < num_components; i++)
892 components_local[i + 1] = xyzw[i + var->data.location_frac];
893
894 components = components_local;
895 }
896 break;
897 default:
898 break;
899 }
900
901 if (var->data.mode & nir_var_system_value) {
902 fprintf(fp, " (%s%s)", loc, components);
903 } else {
904 fprintf(fp, " (%s%s, %u, %u)%s", loc,
905 components,
906 var->data.driver_location, var->data.binding,
907 var->data.compact ? " compact" : "");
908 }
909 }
910
911 if (var->constant_initializer) {
912 if (var->constant_initializer->is_null_constant) {
913 fprintf(fp, " = null");
914 } else {
915 fprintf(fp, " = { ");
916 print_constant(var->constant_initializer, var->type, state);
917 fprintf(fp, " }");
918 }
919 }
920 if (glsl_type_is_sampler(var->type) && var->data.sampler.is_inline_sampler) {
921 fprintf(fp, " = { %s, %s, %s }",
922 get_constant_sampler_addressing_mode(var->data.sampler.addressing_mode),
923 var->data.sampler.normalized_coordinates ? "true" : "false",
924 get_constant_sampler_filter_mode(var->data.sampler.filter_mode));
925 }
926 if (var->pointer_initializer)
927 fprintf(fp, " = &%s", get_var_name(var->pointer_initializer, state));
928
929 fprintf(fp, "\n");
930 print_annotation(state, var);
931 }
932
933 static void
print_deref_link(const nir_deref_instr * instr,bool whole_chain,print_state * state)934 print_deref_link(const nir_deref_instr *instr, bool whole_chain, print_state *state)
935 {
936 FILE *fp = state->fp;
937
938 if (instr->deref_type == nir_deref_type_var) {
939 fprintf(fp, "%s", get_var_name(instr->var, state));
940 return;
941 } else if (instr->deref_type == nir_deref_type_cast) {
942 fprintf(fp, "(%s *)", glsl_get_type_name(instr->type));
943 print_src(&instr->parent, state, nir_type_invalid);
944 return;
945 }
946
947 nir_deref_instr *parent =
948 nir_instr_as_deref(instr->parent.ssa->parent_instr);
949
950 /* Is the parent we're going to print a bare cast? */
951 const bool is_parent_cast =
952 whole_chain && parent->deref_type == nir_deref_type_cast;
953
954 /* If we're not printing the whole chain, the parent we print will be a SSA
955 * value that represents a pointer. The only deref type that naturally
956 * gives a pointer is a cast.
957 */
958 const bool is_parent_pointer =
959 !whole_chain || parent->deref_type == nir_deref_type_cast;
960
961 /* Struct derefs have a nice syntax that works on pointers, arrays derefs
962 * do not.
963 */
964 const bool need_deref =
965 is_parent_pointer && instr->deref_type != nir_deref_type_struct;
966
967 /* Cast need extra parens and so * dereferences */
968 if (is_parent_cast || need_deref)
969 fprintf(fp, "(");
970
971 if (need_deref)
972 fprintf(fp, "*");
973
974 if (whole_chain) {
975 print_deref_link(parent, whole_chain, state);
976 } else {
977 print_src(&instr->parent, state, nir_type_invalid);
978 }
979
980 if (is_parent_cast || need_deref)
981 fprintf(fp, ")");
982
983 switch (instr->deref_type) {
984 case nir_deref_type_struct:
985 fprintf(fp, "%s%s", is_parent_pointer ? "->" : ".",
986 glsl_get_struct_elem_name(parent->type, instr->strct.index));
987 break;
988
989 case nir_deref_type_array:
990 case nir_deref_type_ptr_as_array: {
991 if (nir_src_is_const(instr->arr.index)) {
992 fprintf(fp, "[%" PRId64 "]", nir_src_as_int(instr->arr.index));
993 } else {
994 fprintf(fp, "[");
995 print_src(&instr->arr.index, state, nir_type_invalid);
996 fprintf(fp, "]");
997 }
998 break;
999 }
1000
1001 case nir_deref_type_array_wildcard:
1002 fprintf(fp, "[*]");
1003 break;
1004
1005 default:
1006 unreachable("Invalid deref instruction type");
1007 }
1008 }
1009
1010 static void
print_deref_instr(nir_deref_instr * instr,print_state * state)1011 print_deref_instr(nir_deref_instr *instr, print_state *state)
1012 {
1013 FILE *fp = state->fp;
1014
1015 print_def(&instr->def, state);
1016
1017 switch (instr->deref_type) {
1018 case nir_deref_type_var:
1019 fprintf(fp, " = deref_var ");
1020 break;
1021 case nir_deref_type_array:
1022 case nir_deref_type_array_wildcard:
1023 fprintf(fp, " = deref_array ");
1024 break;
1025 case nir_deref_type_struct:
1026 fprintf(fp, " = deref_struct ");
1027 break;
1028 case nir_deref_type_cast:
1029 fprintf(fp, " = deref_cast ");
1030 break;
1031 case nir_deref_type_ptr_as_array:
1032 fprintf(fp, " = deref_ptr_as_array ");
1033 break;
1034 default:
1035 unreachable("Invalid deref instruction type");
1036 }
1037
1038 /* Only casts naturally return a pointer type */
1039 if (instr->deref_type != nir_deref_type_cast)
1040 fprintf(fp, "&");
1041
1042 print_deref_link(instr, false, state);
1043
1044 fprintf(fp, " (");
1045 unsigned modes = instr->modes;
1046 while (modes) {
1047 int m = u_bit_scan(&modes);
1048 fprintf(fp, "%s%s", get_variable_mode_str(1 << m, true),
1049 modes ? "|" : "");
1050 }
1051 fprintf(fp, " %s)", glsl_get_type_name(instr->type));
1052
1053 if (instr->deref_type == nir_deref_type_cast) {
1054 fprintf(fp, " (ptr_stride=%u, align_mul=%u, align_offset=%u)",
1055 instr->cast.ptr_stride,
1056 instr->cast.align_mul, instr->cast.align_offset);
1057 }
1058
1059 if (instr->deref_type != nir_deref_type_var &&
1060 instr->deref_type != nir_deref_type_cast) {
1061 /* Print the entire chain as a comment */
1062 fprintf(fp, " // &");
1063 print_deref_link(instr, true, state);
1064 }
1065 }
1066
1067 static const char *
vulkan_descriptor_type_name(VkDescriptorType type)1068 vulkan_descriptor_type_name(VkDescriptorType type)
1069 {
1070 switch (type) {
1071 case VK_DESCRIPTOR_TYPE_SAMPLER:
1072 return "sampler";
1073 case VK_DESCRIPTOR_TYPE_COMBINED_IMAGE_SAMPLER:
1074 return "texture+sampler";
1075 case VK_DESCRIPTOR_TYPE_SAMPLED_IMAGE:
1076 return "texture";
1077 case VK_DESCRIPTOR_TYPE_STORAGE_IMAGE:
1078 return "image";
1079 case VK_DESCRIPTOR_TYPE_UNIFORM_TEXEL_BUFFER:
1080 return "texture-buffer";
1081 case VK_DESCRIPTOR_TYPE_STORAGE_TEXEL_BUFFER:
1082 return "image-buffer";
1083 case VK_DESCRIPTOR_TYPE_UNIFORM_BUFFER:
1084 return "UBO";
1085 case VK_DESCRIPTOR_TYPE_STORAGE_BUFFER:
1086 return "SSBO";
1087 case VK_DESCRIPTOR_TYPE_UNIFORM_BUFFER_DYNAMIC:
1088 return "UBO";
1089 case VK_DESCRIPTOR_TYPE_STORAGE_BUFFER_DYNAMIC:
1090 return "SSBO";
1091 case VK_DESCRIPTOR_TYPE_INPUT_ATTACHMENT:
1092 return "input-att";
1093 case VK_DESCRIPTOR_TYPE_INLINE_UNIFORM_BLOCK:
1094 return "inline-UBO";
1095 case VK_DESCRIPTOR_TYPE_ACCELERATION_STRUCTURE_KHR:
1096 return "accel-struct";
1097 default:
1098 return "unknown";
1099 }
1100 }
1101
1102 static void
print_alu_type(nir_alu_type type,print_state * state)1103 print_alu_type(nir_alu_type type, print_state *state)
1104 {
1105 FILE *fp = state->fp;
1106 unsigned size = nir_alu_type_get_type_size(type);
1107 const char *name;
1108
1109 switch (nir_alu_type_get_base_type(type)) {
1110 case nir_type_int:
1111 name = "int";
1112 break;
1113 case nir_type_uint:
1114 name = "uint";
1115 break;
1116 case nir_type_bool:
1117 name = "bool";
1118 break;
1119 case nir_type_float:
1120 name = "float";
1121 break;
1122 default:
1123 name = "invalid";
1124 }
1125 if (size)
1126 fprintf(fp, "%s%u", name, size);
1127 else
1128 fprintf(fp, "%s", name);
1129 }
1130
1131 static void
print_intrinsic_instr(nir_intrinsic_instr * instr,print_state * state)1132 print_intrinsic_instr(nir_intrinsic_instr *instr, print_state *state)
1133 {
1134 const nir_intrinsic_info *info = &nir_intrinsic_infos[instr->intrinsic];
1135 unsigned num_srcs = info->num_srcs;
1136 FILE *fp = state->fp;
1137
1138 if (info->has_dest) {
1139 print_def(&instr->def, state);
1140 fprintf(fp, " = ");
1141 } else {
1142 print_no_dest_padding(state);
1143 }
1144
1145 fprintf(fp, "@%s", info->name);
1146
1147 for (unsigned i = 0; i < num_srcs; i++) {
1148 if (i == 0)
1149 fprintf(fp, " (");
1150 else
1151 fprintf(fp, ", ");
1152
1153 print_src(&instr->src[i], state, nir_intrinsic_instr_src_type(instr, i));
1154 }
1155
1156 if (num_srcs)
1157 fprintf(fp, ")");
1158
1159 for (unsigned i = 0; i < info->num_indices; i++) {
1160 unsigned idx = info->indices[i];
1161 if (i == 0)
1162 fprintf(fp, " (");
1163 else
1164 fprintf(fp, ", ");
1165 switch (idx) {
1166 case NIR_INTRINSIC_WRITE_MASK: {
1167 /* special case wrmask to show it as a writemask.. */
1168 unsigned wrmask = nir_intrinsic_write_mask(instr);
1169 fprintf(fp, "wrmask=");
1170 for (unsigned i = 0; i < instr->num_components; i++)
1171 if ((wrmask >> i) & 1)
1172 fprintf(fp, "%c", comp_mask_string(instr->num_components)[i]);
1173 break;
1174 }
1175
1176 case NIR_INTRINSIC_REDUCTION_OP: {
1177 nir_op reduction_op = nir_intrinsic_reduction_op(instr);
1178 fprintf(fp, "reduction_op=%s", nir_op_infos[reduction_op].name);
1179 break;
1180 }
1181
1182 case NIR_INTRINSIC_ATOMIC_OP: {
1183 nir_atomic_op atomic_op = nir_intrinsic_atomic_op(instr);
1184 fprintf(fp, "atomic_op=");
1185
1186 switch (atomic_op) {
1187 case nir_atomic_op_iadd:
1188 fprintf(fp, "iadd");
1189 break;
1190 case nir_atomic_op_imin:
1191 fprintf(fp, "imin");
1192 break;
1193 case nir_atomic_op_umin:
1194 fprintf(fp, "umin");
1195 break;
1196 case nir_atomic_op_imax:
1197 fprintf(fp, "imax");
1198 break;
1199 case nir_atomic_op_umax:
1200 fprintf(fp, "umax");
1201 break;
1202 case nir_atomic_op_iand:
1203 fprintf(fp, "iand");
1204 break;
1205 case nir_atomic_op_ior:
1206 fprintf(fp, "ior");
1207 break;
1208 case nir_atomic_op_ixor:
1209 fprintf(fp, "ixor");
1210 break;
1211 case nir_atomic_op_xchg:
1212 fprintf(fp, "xchg");
1213 break;
1214 case nir_atomic_op_fadd:
1215 fprintf(fp, "fadd");
1216 break;
1217 case nir_atomic_op_fmin:
1218 fprintf(fp, "fmin");
1219 break;
1220 case nir_atomic_op_fmax:
1221 fprintf(fp, "fmax");
1222 break;
1223 case nir_atomic_op_cmpxchg:
1224 fprintf(fp, "cmpxchg");
1225 break;
1226 case nir_atomic_op_fcmpxchg:
1227 fprintf(fp, "fcmpxchg");
1228 break;
1229 case nir_atomic_op_inc_wrap:
1230 fprintf(fp, "inc_wrap");
1231 break;
1232 case nir_atomic_op_dec_wrap:
1233 fprintf(fp, "dec_wrap");
1234 break;
1235 case nir_atomic_op_ordered_add_gfx12_amd:
1236 fprintf(fp, "ordered_add");
1237 break;
1238 }
1239 break;
1240 }
1241
1242 case NIR_INTRINSIC_IMAGE_DIM: {
1243 static const char *dim_name[] = {
1244 [GLSL_SAMPLER_DIM_1D] = "1D",
1245 [GLSL_SAMPLER_DIM_2D] = "2D",
1246 [GLSL_SAMPLER_DIM_3D] = "3D",
1247 [GLSL_SAMPLER_DIM_CUBE] = "Cube",
1248 [GLSL_SAMPLER_DIM_RECT] = "Rect",
1249 [GLSL_SAMPLER_DIM_BUF] = "Buf",
1250 [GLSL_SAMPLER_DIM_MS] = "2D-MSAA",
1251 [GLSL_SAMPLER_DIM_SUBPASS] = "Subpass",
1252 [GLSL_SAMPLER_DIM_SUBPASS_MS] = "Subpass-MSAA",
1253 };
1254 enum glsl_sampler_dim dim = nir_intrinsic_image_dim(instr);
1255 assert(dim < ARRAY_SIZE(dim_name) && dim_name[dim]);
1256 fprintf(fp, "image_dim=%s", dim_name[dim]);
1257 break;
1258 }
1259
1260 case NIR_INTRINSIC_IMAGE_ARRAY: {
1261 bool array = nir_intrinsic_image_array(instr);
1262 fprintf(fp, "image_array=%s", array ? "true" : "false");
1263 break;
1264 }
1265
1266 case NIR_INTRINSIC_FORMAT: {
1267 enum pipe_format format = nir_intrinsic_format(instr);
1268 fprintf(fp, "format=%s", util_format_short_name(format));
1269 break;
1270 }
1271
1272 case NIR_INTRINSIC_DESC_TYPE: {
1273 VkDescriptorType desc_type = nir_intrinsic_desc_type(instr);
1274 fprintf(fp, "desc_type=%s", vulkan_descriptor_type_name(desc_type));
1275 break;
1276 }
1277
1278 case NIR_INTRINSIC_SRC_TYPE: {
1279 fprintf(fp, "src_type=");
1280 print_alu_type(nir_intrinsic_src_type(instr), state);
1281 break;
1282 }
1283
1284 case NIR_INTRINSIC_DEST_TYPE: {
1285 fprintf(fp, "dest_type=");
1286 print_alu_type(nir_intrinsic_dest_type(instr), state);
1287 break;
1288 }
1289
1290 case NIR_INTRINSIC_SWIZZLE_MASK: {
1291 fprintf(fp, "swizzle_mask=");
1292 unsigned mask = nir_intrinsic_swizzle_mask(instr);
1293 if (instr->intrinsic == nir_intrinsic_quad_swizzle_amd) {
1294 for (unsigned i = 0; i < 4; i++)
1295 fprintf(fp, "%d", (mask >> (i * 2) & 3));
1296 } else if (instr->intrinsic == nir_intrinsic_masked_swizzle_amd) {
1297 fprintf(fp, "((id & %d) | %d) ^ %d", mask & 0x1F,
1298 (mask >> 5) & 0x1F,
1299 (mask >> 10) & 0x1F);
1300 } else {
1301 fprintf(fp, "%d", mask);
1302 }
1303 break;
1304 }
1305
1306 case NIR_INTRINSIC_MEMORY_SEMANTICS: {
1307 nir_memory_semantics semantics = nir_intrinsic_memory_semantics(instr);
1308 fprintf(fp, "mem_semantics=");
1309 switch (semantics & (NIR_MEMORY_ACQUIRE | NIR_MEMORY_RELEASE)) {
1310 case 0:
1311 fprintf(fp, "NONE");
1312 break;
1313 case NIR_MEMORY_ACQUIRE:
1314 fprintf(fp, "ACQ");
1315 break;
1316 case NIR_MEMORY_RELEASE:
1317 fprintf(fp, "REL");
1318 break;
1319 default:
1320 fprintf(fp, "ACQ|REL");
1321 break;
1322 }
1323 if (semantics & (NIR_MEMORY_MAKE_AVAILABLE))
1324 fprintf(fp, "|AVAILABLE");
1325 if (semantics & (NIR_MEMORY_MAKE_VISIBLE))
1326 fprintf(fp, "|VISIBLE");
1327 break;
1328 }
1329
1330 case NIR_INTRINSIC_MEMORY_MODES: {
1331 fprintf(fp, "mem_modes=");
1332 unsigned int modes = nir_intrinsic_memory_modes(instr);
1333 if (modes == 0)
1334 fputc('0', fp);
1335 while (modes) {
1336 nir_variable_mode m = u_bit_scan(&modes);
1337 fprintf(fp, "%s%s", get_variable_mode_str(1 << m, true), modes ? "|" : "");
1338 }
1339 break;
1340 }
1341
1342 case NIR_INTRINSIC_EXECUTION_SCOPE:
1343 case NIR_INTRINSIC_MEMORY_SCOPE: {
1344 mesa_scope scope =
1345 idx == NIR_INTRINSIC_MEMORY_SCOPE ? nir_intrinsic_memory_scope(instr)
1346 : nir_intrinsic_execution_scope(instr);
1347 const char *name = mesa_scope_name(scope);
1348 static const char prefix[] = "SCOPE_";
1349 if (strncmp(name, prefix, sizeof(prefix) - 1) == 0)
1350 name += sizeof(prefix) - 1;
1351 fprintf(fp, "%s=%s", nir_intrinsic_index_names[idx], name);
1352 break;
1353 }
1354
1355 case NIR_INTRINSIC_IO_SEMANTICS: {
1356 struct nir_io_semantics io = nir_intrinsic_io_semantics(instr);
1357
1358 /* Try to figure out the mode so we can interpret the location */
1359 nir_variable_mode mode = nir_var_mem_generic;
1360 switch (instr->intrinsic) {
1361 case nir_intrinsic_load_input:
1362 case nir_intrinsic_load_per_primitive_input:
1363 case nir_intrinsic_load_interpolated_input:
1364 case nir_intrinsic_load_per_vertex_input:
1365 case nir_intrinsic_load_input_vertex:
1366 case nir_intrinsic_load_coefficients_agx:
1367 case nir_intrinsic_load_attribute_pan:
1368 mode = nir_var_shader_in;
1369 break;
1370
1371 case nir_intrinsic_load_output:
1372 case nir_intrinsic_store_output:
1373 case nir_intrinsic_store_per_primitive_output:
1374 case nir_intrinsic_store_per_vertex_output:
1375 case nir_intrinsic_store_per_view_output:
1376 mode = nir_var_shader_out;
1377 break;
1378
1379 default:
1380 break;
1381 }
1382
1383 /* Using that mode, we should be able to name the location */
1384 char buf[4];
1385 const char *loc = get_location_str(io.location,
1386 state->shader->info.stage, mode,
1387 buf);
1388
1389 fprintf(fp, "io location=%s slots=%u", loc, io.num_slots);
1390
1391 if (io.interp_explicit_strict)
1392 fprintf(fp, " explicit_strict");
1393
1394 if (io.dual_source_blend_index)
1395 fprintf(fp, " dualsrc");
1396
1397 if (io.fb_fetch_output)
1398 fprintf(fp, " fbfetch");
1399
1400 if (io.fb_fetch_output_coherent)
1401 fprintf(fp, " coherent");
1402
1403 if (io.per_view)
1404 fprintf(fp, " perview");
1405
1406 if (io.medium_precision)
1407 fprintf(fp, " mediump");
1408
1409 if (io.high_16bits)
1410 fprintf(fp, " high_16bits");
1411
1412 if (io.invariant)
1413 fprintf(fp, " invariant");
1414
1415 if (io.high_dvec2)
1416 fprintf(fp, " high_dvec2");
1417
1418 if (io.no_varying)
1419 fprintf(fp, " no_varying");
1420
1421 if (io.no_sysval_output)
1422 fprintf(fp, " no_sysval_output");
1423
1424 if (state->shader &&
1425 state->shader->info.stage == MESA_SHADER_GEOMETRY &&
1426 (instr->intrinsic == nir_intrinsic_store_output ||
1427 instr->intrinsic == nir_intrinsic_store_per_primitive_output ||
1428 instr->intrinsic == nir_intrinsic_store_per_vertex_output ||
1429 instr->intrinsic == nir_intrinsic_store_per_view_output)) {
1430 unsigned gs_streams = io.gs_streams;
1431 fprintf(fp, " gs_streams(");
1432 for (unsigned i = 0; i < 4; i++) {
1433 fprintf(fp, "%s%c=%u", i ? " " : "", "xyzw"[i],
1434 (gs_streams >> (i * 2)) & 0x3);
1435 }
1436 fprintf(fp, ")");
1437 }
1438
1439 break;
1440 }
1441
1442 case NIR_INTRINSIC_IO_XFB:
1443 case NIR_INTRINSIC_IO_XFB2: {
1444 /* This prints both IO_XFB and IO_XFB2. */
1445 fprintf(fp, "xfb%s(", idx == NIR_INTRINSIC_IO_XFB ? "" : "2");
1446 bool first = true;
1447 for (unsigned i = 0; i < 2; i++) {
1448 unsigned start_comp = (idx == NIR_INTRINSIC_IO_XFB ? 0 : 2) + i;
1449 nir_io_xfb xfb = start_comp < 2 ? nir_intrinsic_io_xfb(instr) : nir_intrinsic_io_xfb2(instr);
1450
1451 if (!xfb.out[i].num_components)
1452 continue;
1453
1454 if (!first)
1455 fprintf(fp, ", ");
1456 first = false;
1457
1458 if (xfb.out[i].num_components > 1) {
1459 fprintf(fp, "components=%u..%u",
1460 start_comp, start_comp + xfb.out[i].num_components - 1);
1461 } else {
1462 fprintf(fp, "component=%u", start_comp);
1463 }
1464 fprintf(fp, " buffer=%u offset=%u",
1465 xfb.out[i].buffer, (uint32_t)xfb.out[i].offset * 4);
1466 }
1467 fprintf(fp, ")");
1468 break;
1469 }
1470
1471 case NIR_INTRINSIC_ROUNDING_MODE: {
1472 fprintf(fp, "rounding_mode=");
1473 switch (nir_intrinsic_rounding_mode(instr)) {
1474 case nir_rounding_mode_undef:
1475 fprintf(fp, "undef");
1476 break;
1477 case nir_rounding_mode_rtne:
1478 fprintf(fp, "rtne");
1479 break;
1480 case nir_rounding_mode_ru:
1481 fprintf(fp, "ru");
1482 break;
1483 case nir_rounding_mode_rd:
1484 fprintf(fp, "rd");
1485 break;
1486 case nir_rounding_mode_rtz:
1487 fprintf(fp, "rtz");
1488 break;
1489 default:
1490 fprintf(fp, "unknown");
1491 break;
1492 }
1493 break;
1494 }
1495
1496 case NIR_INTRINSIC_RAY_QUERY_VALUE: {
1497 fprintf(fp, "ray_query_value=");
1498 switch (nir_intrinsic_ray_query_value(instr)) {
1499 #define VAL(_name) \
1500 case nir_ray_query_value_##_name: \
1501 fprintf(fp, #_name); \
1502 break
1503 VAL(intersection_type);
1504 VAL(intersection_t);
1505 VAL(intersection_instance_custom_index);
1506 VAL(intersection_instance_id);
1507 VAL(intersection_instance_sbt_index);
1508 VAL(intersection_geometry_index);
1509 VAL(intersection_primitive_index);
1510 VAL(intersection_barycentrics);
1511 VAL(intersection_front_face);
1512 VAL(intersection_object_ray_direction);
1513 VAL(intersection_object_ray_origin);
1514 VAL(intersection_object_to_world);
1515 VAL(intersection_world_to_object);
1516 VAL(intersection_candidate_aabb_opaque);
1517 VAL(tmin);
1518 VAL(flags);
1519 VAL(world_ray_direction);
1520 VAL(world_ray_origin);
1521 #undef VAL
1522 default:
1523 fprintf(fp, "unknown");
1524 break;
1525 }
1526 break;
1527 }
1528
1529 case NIR_INTRINSIC_RESOURCE_ACCESS_INTEL: {
1530 fprintf(fp, "resource_intel=");
1531 unsigned int modes = nir_intrinsic_resource_access_intel(instr);
1532 if (modes == 0)
1533 fputc('0', fp);
1534 while (modes) {
1535 nir_resource_data_intel i = 1u << u_bit_scan(&modes);
1536 switch (i) {
1537 case nir_resource_intel_bindless:
1538 fprintf(fp, "bindless");
1539 break;
1540 case nir_resource_intel_pushable:
1541 fprintf(fp, "pushable");
1542 break;
1543 case nir_resource_intel_sampler:
1544 fprintf(fp, "sampler");
1545 break;
1546 case nir_resource_intel_non_uniform:
1547 fprintf(fp, "non-uniform");
1548 break;
1549 case nir_resource_intel_sampler_embedded:
1550 fprintf(fp, "sampler-embedded");
1551 break;
1552 default:
1553 fprintf(fp, "unknown");
1554 break;
1555 }
1556 fprintf(fp, "%s", modes ? "|" : "");
1557 }
1558 break;
1559 }
1560
1561 case NIR_INTRINSIC_ACCESS: {
1562 fprintf(fp, "access=");
1563 print_access(nir_intrinsic_access(instr), state, "|");
1564 break;
1565 }
1566
1567 case NIR_INTRINSIC_MATRIX_LAYOUT: {
1568 fprintf(fp, "matrix_layout=");
1569 switch (nir_intrinsic_matrix_layout(instr)) {
1570 case GLSL_MATRIX_LAYOUT_ROW_MAJOR:
1571 fprintf(fp, "row_major");
1572 break;
1573 case GLSL_MATRIX_LAYOUT_COLUMN_MAJOR:
1574 fprintf(fp, "col_major");
1575 break;
1576 default:
1577 fprintf(fp, "unknown");
1578 break;
1579 }
1580 break;
1581 }
1582
1583 case NIR_INTRINSIC_CMAT_DESC: {
1584 struct glsl_cmat_description desc = nir_intrinsic_cmat_desc(instr);
1585 const struct glsl_type *t = glsl_cmat_type(&desc);
1586 fprintf(fp, "%s", glsl_get_type_name(t));
1587 break;
1588 }
1589
1590 case NIR_INTRINSIC_CMAT_SIGNED_MASK: {
1591 fprintf(fp, "cmat_signed=");
1592 unsigned int mask = nir_intrinsic_cmat_signed_mask(instr);
1593 if (mask == 0)
1594 fputc('0', fp);
1595 while (mask) {
1596 nir_cmat_signed i = 1u << u_bit_scan(&mask);
1597 switch (i) {
1598 case NIR_CMAT_A_SIGNED:
1599 fputc('A', fp);
1600 break;
1601 case NIR_CMAT_B_SIGNED:
1602 fputc('B', fp);
1603 break;
1604 case NIR_CMAT_C_SIGNED:
1605 fputc('C', fp);
1606 break;
1607 case NIR_CMAT_RESULT_SIGNED:
1608 fprintf(fp, "Result");
1609 break;
1610 default:
1611 fprintf(fp, "unknown");
1612 break;
1613 }
1614 fprintf(fp, "%s", mask ? "|" : "");
1615 }
1616 break;
1617 }
1618
1619 case NIR_INTRINSIC_ALU_OP: {
1620 nir_op alu_op = nir_intrinsic_alu_op(instr);
1621 fprintf(fp, "alu_op=%s", nir_op_infos[alu_op].name);
1622 break;
1623 }
1624
1625 case NIR_INTRINSIC_INTERP_MODE:
1626 fprintf(fp, "interp_mode=%s",
1627 glsl_interp_mode_name(nir_intrinsic_interp_mode(instr)));
1628 break;
1629
1630 default: {
1631 unsigned off = info->index_map[idx] - 1;
1632 fprintf(fp, "%s=%d", nir_intrinsic_index_names[idx], instr->const_index[off]);
1633 break;
1634 }
1635 }
1636 }
1637 if (info->num_indices)
1638 fprintf(fp, ")");
1639
1640 if (!state->shader)
1641 return;
1642
1643 nir_variable_mode var_mode;
1644 switch (instr->intrinsic) {
1645 case nir_intrinsic_load_uniform:
1646 var_mode = nir_var_uniform;
1647 break;
1648 case nir_intrinsic_load_input:
1649 case nir_intrinsic_load_per_primitive_input:
1650 case nir_intrinsic_load_interpolated_input:
1651 case nir_intrinsic_load_per_vertex_input:
1652 var_mode = nir_var_shader_in;
1653 break;
1654 case nir_intrinsic_load_output:
1655 case nir_intrinsic_store_output:
1656 case nir_intrinsic_store_per_vertex_output:
1657 case nir_intrinsic_store_per_view_output:
1658 var_mode = nir_var_shader_out;
1659 break;
1660 default:
1661 return;
1662 }
1663
1664 if (instr->name) {
1665 fprintf(fp, " // %s", instr->name);
1666 return;
1667 }
1668
1669 nir_foreach_variable_with_modes(var, state->shader, var_mode) {
1670 if (!var->name)
1671 continue;
1672
1673 bool match;
1674 if (instr->intrinsic == nir_intrinsic_load_uniform) {
1675 match = var->data.driver_location == nir_intrinsic_base(instr);
1676 } else {
1677 match = nir_intrinsic_component(instr) >= var->data.location_frac &&
1678 nir_intrinsic_component(instr) <
1679 (var->data.location_frac + glsl_get_components(var->type));
1680 }
1681
1682 if (match) {
1683 fprintf(fp, " // %s", var->name);
1684 break;
1685 }
1686 }
1687 }
1688
1689 static void
print_tex_instr(nir_tex_instr * instr,print_state * state)1690 print_tex_instr(nir_tex_instr *instr, print_state *state)
1691 {
1692 FILE *fp = state->fp;
1693
1694 print_def(&instr->def, state);
1695
1696 fprintf(fp, " = (");
1697 print_alu_type(instr->dest_type, state);
1698 fprintf(fp, ")");
1699
1700 switch (instr->op) {
1701 case nir_texop_tex:
1702 fprintf(fp, "tex ");
1703 break;
1704 case nir_texop_txb:
1705 fprintf(fp, "txb ");
1706 break;
1707 case nir_texop_txl:
1708 fprintf(fp, "txl ");
1709 break;
1710 case nir_texop_txd:
1711 fprintf(fp, "txd ");
1712 break;
1713 case nir_texop_txf:
1714 fprintf(fp, "txf ");
1715 break;
1716 case nir_texop_txf_ms:
1717 fprintf(fp, "txf_ms ");
1718 break;
1719 case nir_texop_txf_ms_fb:
1720 fprintf(fp, "txf_ms_fb ");
1721 break;
1722 case nir_texop_txf_ms_mcs_intel:
1723 fprintf(fp, "txf_ms_mcs_intel ");
1724 break;
1725 case nir_texop_txs:
1726 fprintf(fp, "txs ");
1727 break;
1728 case nir_texop_lod:
1729 fprintf(fp, "lod ");
1730 break;
1731 case nir_texop_tg4:
1732 fprintf(fp, "tg4 ");
1733 break;
1734 case nir_texop_query_levels:
1735 fprintf(fp, "query_levels ");
1736 break;
1737 case nir_texop_texture_samples:
1738 fprintf(fp, "texture_samples ");
1739 break;
1740 case nir_texop_samples_identical:
1741 fprintf(fp, "samples_identical ");
1742 break;
1743 case nir_texop_tex_prefetch:
1744 fprintf(fp, "tex (pre-dispatchable) ");
1745 break;
1746 case nir_texop_fragment_fetch_amd:
1747 fprintf(fp, "fragment_fetch_amd ");
1748 break;
1749 case nir_texop_fragment_mask_fetch_amd:
1750 fprintf(fp, "fragment_mask_fetch_amd ");
1751 break;
1752 case nir_texop_descriptor_amd:
1753 fprintf(fp, "descriptor_amd ");
1754 break;
1755 case nir_texop_sampler_descriptor_amd:
1756 fprintf(fp, "sampler_descriptor_amd ");
1757 break;
1758 case nir_texop_lod_bias_agx:
1759 fprintf(fp, "lod_bias_agx ");
1760 break;
1761 case nir_texop_has_custom_border_color_agx:
1762 fprintf(fp, "has_custom_border_color_agx ");
1763 break;
1764 case nir_texop_custom_border_color_agx:
1765 fprintf(fp, "custom_border_color_agx ");
1766 break;
1767 case nir_texop_hdr_dim_nv:
1768 fprintf(fp, "hdr_dim_nv ");
1769 break;
1770 case nir_texop_tex_type_nv:
1771 fprintf(fp, "tex_type_nv ");
1772 break;
1773 default:
1774 unreachable("Invalid texture operation");
1775 break;
1776 }
1777
1778 bool has_texture_deref = false, has_sampler_deref = false;
1779 for (unsigned i = 0; i < instr->num_srcs; i++) {
1780 if (i > 0) {
1781 fprintf(fp, ", ");
1782 }
1783
1784 print_src(&instr->src[i].src, state, nir_tex_instr_src_type(instr, i));
1785 fprintf(fp, " ");
1786
1787 switch (instr->src[i].src_type) {
1788 case nir_tex_src_backend1:
1789 fprintf(fp, "(backend1)");
1790 break;
1791 case nir_tex_src_backend2:
1792 fprintf(fp, "(backend2)");
1793 break;
1794 case nir_tex_src_coord:
1795 fprintf(fp, "(coord)");
1796 break;
1797 case nir_tex_src_projector:
1798 fprintf(fp, "(projector)");
1799 break;
1800 case nir_tex_src_comparator:
1801 fprintf(fp, "(comparator)");
1802 break;
1803 case nir_tex_src_offset:
1804 fprintf(fp, "(offset)");
1805 break;
1806 case nir_tex_src_bias:
1807 fprintf(fp, "(bias)");
1808 break;
1809 case nir_tex_src_lod:
1810 fprintf(fp, "(lod)");
1811 break;
1812 case nir_tex_src_min_lod:
1813 fprintf(fp, "(min_lod)");
1814 break;
1815 case nir_tex_src_ms_index:
1816 fprintf(fp, "(ms_index)");
1817 break;
1818 case nir_tex_src_ms_mcs_intel:
1819 fprintf(fp, "(ms_mcs_intel)");
1820 break;
1821 case nir_tex_src_ddx:
1822 fprintf(fp, "(ddx)");
1823 break;
1824 case nir_tex_src_ddy:
1825 fprintf(fp, "(ddy)");
1826 break;
1827 case nir_tex_src_sampler_deref_intrinsic:
1828 has_texture_deref = true;
1829 fprintf(fp, "(sampler_deref_intrinsic)");
1830 break;
1831 case nir_tex_src_texture_deref_intrinsic:
1832 has_texture_deref = true;
1833 fprintf(fp, "(texture_deref_intrinsic)");
1834 break;
1835 case nir_tex_src_texture_deref:
1836 has_texture_deref = true;
1837 fprintf(fp, "(texture_deref)");
1838 break;
1839 case nir_tex_src_sampler_deref:
1840 has_sampler_deref = true;
1841 fprintf(fp, "(sampler_deref)");
1842 break;
1843 case nir_tex_src_texture_offset:
1844 fprintf(fp, "(texture_offset)");
1845 break;
1846 case nir_tex_src_sampler_offset:
1847 fprintf(fp, "(sampler_offset)");
1848 break;
1849 case nir_tex_src_texture_handle:
1850 fprintf(fp, "(texture_handle)");
1851 break;
1852 case nir_tex_src_sampler_handle:
1853 fprintf(fp, "(sampler_handle)");
1854 break;
1855 case nir_tex_src_plane:
1856 fprintf(fp, "(plane)");
1857 break;
1858
1859 default:
1860 unreachable("Invalid texture source type");
1861 break;
1862 }
1863 }
1864
1865 if (instr->is_gather_implicit_lod)
1866 fprintf(fp, ", implicit lod");
1867
1868 if (instr->op == nir_texop_tg4) {
1869 fprintf(fp, ", %u (gather_component)", instr->component);
1870 }
1871
1872 if (nir_tex_instr_has_explicit_tg4_offsets(instr)) {
1873 fprintf(fp, ", { (%i, %i)", instr->tg4_offsets[0][0], instr->tg4_offsets[0][1]);
1874 for (unsigned i = 1; i < 4; ++i)
1875 fprintf(fp, ", (%i, %i)", instr->tg4_offsets[i][0],
1876 instr->tg4_offsets[i][1]);
1877 fprintf(fp, " } (offsets)");
1878 }
1879
1880 if (instr->op != nir_texop_txf_ms_fb && !has_texture_deref) {
1881 fprintf(fp, ", %u (texture)", instr->texture_index);
1882 }
1883
1884 if (nir_tex_instr_need_sampler(instr) && !has_sampler_deref) {
1885 fprintf(fp, ", %u (sampler)", instr->sampler_index);
1886 }
1887
1888 if (instr->texture_non_uniform) {
1889 fprintf(fp, ", texture non-uniform");
1890 }
1891
1892 if (instr->sampler_non_uniform) {
1893 fprintf(fp, ", sampler non-uniform");
1894 }
1895
1896 if (instr->is_sparse) {
1897 fprintf(fp, ", sparse");
1898 }
1899 }
1900
1901 static void
print_call_instr(nir_call_instr * instr,print_state * state)1902 print_call_instr(nir_call_instr *instr, print_state *state)
1903 {
1904 FILE *fp = state->fp;
1905
1906 print_no_dest_padding(state);
1907
1908 bool indirect = instr->indirect_callee.ssa;
1909
1910 fprintf(fp, "call %s ", instr->callee->name);
1911 if (indirect) {
1912 fprintf(fp, "(indirect ");
1913 print_src(&instr->indirect_callee, state, nir_type_invalid);
1914 fprintf(fp, ") ");
1915 }
1916
1917 for (unsigned i = 0; i < instr->num_params; i++) {
1918 if (i != 0)
1919 fprintf(fp, ", ");
1920
1921 if (instr->callee->params[i].name)
1922 fprintf(fp, "%s ", instr->callee->params[i].name);
1923
1924 print_src(&instr->params[i], state, nir_type_invalid);
1925 }
1926 }
1927
1928 static void
print_jump_instr(nir_jump_instr * instr,print_state * state)1929 print_jump_instr(nir_jump_instr *instr, print_state *state)
1930 {
1931 FILE *fp = state->fp;
1932
1933 print_no_dest_padding(state);
1934
1935 switch (instr->type) {
1936 case nir_jump_break:
1937 fprintf(fp, "break");
1938 break;
1939
1940 case nir_jump_continue:
1941 fprintf(fp, "continue");
1942 break;
1943
1944 case nir_jump_return:
1945 fprintf(fp, "return");
1946 break;
1947
1948 case nir_jump_halt:
1949 fprintf(fp, "halt");
1950 break;
1951
1952 case nir_jump_goto:
1953 fprintf(fp, "goto b%u",
1954 instr->target ? instr->target->index : -1);
1955 break;
1956
1957 case nir_jump_goto_if:
1958 fprintf(fp, "goto b%u if ",
1959 instr->target ? instr->target->index : -1);
1960 print_src(&instr->condition, state, nir_type_invalid);
1961 fprintf(fp, " else b%u",
1962 instr->else_target ? instr->else_target->index : -1);
1963 break;
1964 }
1965 }
1966
1967 static void
print_ssa_undef_instr(nir_undef_instr * instr,print_state * state)1968 print_ssa_undef_instr(nir_undef_instr *instr, print_state *state)
1969 {
1970 FILE *fp = state->fp;
1971 print_def(&instr->def, state);
1972 fprintf(fp, " = undefined");
1973 }
1974
1975 static void
print_phi_instr(nir_phi_instr * instr,print_state * state)1976 print_phi_instr(nir_phi_instr *instr, print_state *state)
1977 {
1978 FILE *fp = state->fp;
1979 print_def(&instr->def, state);
1980 fprintf(fp, " = phi ");
1981 nir_foreach_phi_src(src, instr) {
1982 if (&src->node != exec_list_get_head(&instr->srcs))
1983 fprintf(fp, ", ");
1984
1985 fprintf(fp, "b%u: ", src->pred->index);
1986 print_src(&src->src, state, nir_type_invalid);
1987 }
1988 }
1989
1990 static void
print_parallel_copy_instr(nir_parallel_copy_instr * instr,print_state * state)1991 print_parallel_copy_instr(nir_parallel_copy_instr *instr, print_state *state)
1992 {
1993 FILE *fp = state->fp;
1994 nir_foreach_parallel_copy_entry(entry, instr) {
1995 if (&entry->node != exec_list_get_head(&instr->entries))
1996 fprintf(fp, "; ");
1997
1998 if (entry->dest_is_reg) {
1999 fprintf(fp, "*");
2000 print_src(&entry->dest.reg, state, nir_type_invalid);
2001 } else {
2002 print_def(&entry->dest.def, state);
2003 }
2004 fprintf(fp, " = ");
2005
2006 if (entry->src_is_reg)
2007 fprintf(fp, "*");
2008 print_src(&entry->src, state, nir_type_invalid);
2009 }
2010 }
2011
2012 static void
print_debug_info_instr(nir_debug_info_instr * instr,print_state * state)2013 print_debug_info_instr(nir_debug_info_instr *instr, print_state *state)
2014 {
2015 FILE *fp = state->fp;
2016
2017 switch (instr->type) {
2018 case nir_debug_info_src_loc:
2019 fprintf(fp, "// 0x%x", instr->src_loc.spirv_offset);
2020 if (instr->src_loc.line)
2021 fprintf(fp, " %s:%u:%u", nir_src_as_string(instr->src_loc.filename), instr->src_loc.line, instr->src_loc.column);
2022 return;
2023 case nir_debug_info_string:
2024 return; /* Strings are printed for their uses. */
2025 }
2026
2027 unreachable("Unimplemented nir_debug_info_type");
2028 }
2029
2030 static void
print_instr(const nir_instr * instr,print_state * state,unsigned tabs)2031 print_instr(const nir_instr *instr, print_state *state, unsigned tabs)
2032 {
2033 FILE *fp = state->fp;
2034
2035 if (state->debug_info) {
2036 nir_debug_info_instr *di = state->debug_info[instr->index];
2037 if (di)
2038 di->src_loc.column = (uint32_t)ftell(fp);
2039 }
2040
2041 print_indentation(tabs, fp);
2042
2043 switch (instr->type) {
2044 case nir_instr_type_alu:
2045 print_alu_instr(nir_instr_as_alu(instr), state);
2046 break;
2047
2048 case nir_instr_type_deref:
2049 print_deref_instr(nir_instr_as_deref(instr), state);
2050 break;
2051
2052 case nir_instr_type_call:
2053 print_call_instr(nir_instr_as_call(instr), state);
2054 break;
2055
2056 case nir_instr_type_intrinsic:
2057 print_intrinsic_instr(nir_instr_as_intrinsic(instr), state);
2058 break;
2059
2060 case nir_instr_type_tex:
2061 print_tex_instr(nir_instr_as_tex(instr), state);
2062 break;
2063
2064 case nir_instr_type_load_const:
2065 print_load_const_instr(nir_instr_as_load_const(instr), state);
2066 break;
2067
2068 case nir_instr_type_jump:
2069 print_jump_instr(nir_instr_as_jump(instr), state);
2070 break;
2071
2072 case nir_instr_type_undef:
2073 print_ssa_undef_instr(nir_instr_as_undef(instr), state);
2074 break;
2075
2076 case nir_instr_type_phi:
2077 print_phi_instr(nir_instr_as_phi(instr), state);
2078 break;
2079
2080 case nir_instr_type_parallel_copy:
2081 print_parallel_copy_instr(nir_instr_as_parallel_copy(instr), state);
2082 break;
2083
2084 case nir_instr_type_debug_info:
2085 print_debug_info_instr(nir_instr_as_debug_info(instr), state);
2086 break;
2087
2088 default:
2089 unreachable("Invalid instruction type");
2090 break;
2091 }
2092
2093 if (NIR_DEBUG(PRINT_PASS_FLAGS) && instr->pass_flags)
2094 fprintf(fp, " (pass_flags: 0x%x)", instr->pass_flags);
2095 }
2096
2097 static bool
block_has_instruction_with_dest(nir_block * block)2098 block_has_instruction_with_dest(nir_block *block)
2099 {
2100 nir_foreach_instr(instr, block) {
2101 switch (instr->type) {
2102 case nir_instr_type_load_const:
2103 case nir_instr_type_deref:
2104 case nir_instr_type_alu:
2105 case nir_instr_type_tex:
2106 case nir_instr_type_undef:
2107 case nir_instr_type_phi:
2108 case nir_instr_type_parallel_copy:
2109 return true;
2110
2111 case nir_instr_type_intrinsic: {
2112 nir_intrinsic_instr *intrin = nir_instr_as_intrinsic(instr);
2113 const nir_intrinsic_info *info = &nir_intrinsic_infos[intrin->intrinsic];
2114 if (info->has_dest)
2115 return true;
2116
2117 /* Doesn't define a new value. */
2118 break;
2119 }
2120
2121 case nir_instr_type_jump:
2122 case nir_instr_type_call:
2123 case nir_instr_type_debug_info:
2124 /* Doesn't define a new value. */
2125 break;
2126 }
2127 }
2128
2129 return false;
2130 }
2131
2132 static void print_cf_node(nir_cf_node *node, print_state *state,
2133 unsigned tabs);
2134
2135 static void
print_block_preds(nir_block * block,print_state * state)2136 print_block_preds(nir_block *block, print_state *state)
2137 {
2138 FILE *fp = state->fp;
2139 nir_block **preds = nir_block_get_predecessors_sorted(block, NULL);
2140 for (unsigned i = 0; i < block->predecessors->entries; i++) {
2141 fprintf(fp, " b%u", preds[i]->index);
2142 }
2143 ralloc_free(preds);
2144 }
2145
2146 static void
print_block_succs(nir_block * block,print_state * state)2147 print_block_succs(nir_block *block, print_state *state)
2148 {
2149 FILE *fp = state->fp;
2150 for (unsigned i = 0; i < 2; i++) {
2151 if (block->successors[i]) {
2152 fprintf(fp, " b%u", block->successors[i]->index);
2153 }
2154 }
2155 }
2156
2157 static void
print_block(nir_block * block,print_state * state,unsigned tabs)2158 print_block(nir_block *block, print_state *state, unsigned tabs)
2159 {
2160 FILE *fp = state->fp;
2161
2162 if (block_has_instruction_with_dest(block))
2163 state->padding_for_no_dest = calculate_padding_for_no_dest(state);
2164 else
2165 state->padding_for_no_dest = 0;
2166
2167 print_indentation(tabs, fp);
2168 fprintf(fp, "%sblock b%u:",
2169 divergence_status(state, block->divergent),
2170 block->index);
2171
2172 const bool empty_block = exec_list_is_empty(&block->instr_list);
2173 if (empty_block) {
2174 fprintf(fp, " // preds:");
2175 print_block_preds(block, state);
2176 fprintf(fp, ", succs:");
2177 print_block_succs(block, state);
2178 fprintf(fp, "\n");
2179 return;
2180 }
2181
2182 const unsigned block_length = 7 + count_digits(block->index) + 1;
2183 const unsigned pred_padding = block_length < state->padding_for_no_dest ? state->padding_for_no_dest - block_length : 0;
2184
2185 fprintf(fp, "%*s// preds:", pred_padding, "");
2186 print_block_preds(block, state);
2187 fprintf(fp, "\n");
2188
2189 nir_foreach_instr(instr, block) {
2190 print_instr(instr, state, tabs);
2191 fprintf(fp, "\n");
2192 print_annotation(state, instr);
2193 }
2194
2195 print_indentation(tabs, fp);
2196 fprintf(fp, "%*s// succs:", state->padding_for_no_dest, "");
2197 print_block_succs(block, state);
2198 fprintf(fp, "\n");
2199 }
2200
2201 static void
print_if(nir_if * if_stmt,print_state * state,unsigned tabs)2202 print_if(nir_if *if_stmt, print_state *state, unsigned tabs)
2203 {
2204 FILE *fp = state->fp;
2205
2206 print_indentation(tabs, fp);
2207 fprintf(fp, "if ");
2208 print_src(&if_stmt->condition, state, nir_type_invalid);
2209 switch (if_stmt->control) {
2210 case nir_selection_control_flatten:
2211 fprintf(fp, " // flatten");
2212 break;
2213 case nir_selection_control_dont_flatten:
2214 fprintf(fp, " // don't flatten");
2215 break;
2216 case nir_selection_control_divergent_always_taken:
2217 fprintf(fp, " // divergent always taken");
2218 break;
2219 case nir_selection_control_none:
2220 default:
2221 break;
2222 }
2223 fprintf(fp, " {\n");
2224 foreach_list_typed(nir_cf_node, node, node, &if_stmt->then_list) {
2225 print_cf_node(node, state, tabs + 1);
2226 }
2227 print_indentation(tabs, fp);
2228 fprintf(fp, "} else {\n");
2229 foreach_list_typed(nir_cf_node, node, node, &if_stmt->else_list) {
2230 print_cf_node(node, state, tabs + 1);
2231 }
2232 print_indentation(tabs, fp);
2233 fprintf(fp, "}\n");
2234 }
2235
2236 static void
print_loop(nir_loop * loop,print_state * state,unsigned tabs)2237 print_loop(nir_loop *loop, print_state *state, unsigned tabs)
2238 {
2239 FILE *fp = state->fp;
2240
2241 print_indentation(tabs, fp);
2242 fprintf(fp, "%sloop {\n", divergence_status(state, loop->divergent_break));
2243 foreach_list_typed(nir_cf_node, node, node, &loop->body) {
2244 print_cf_node(node, state, tabs + 1);
2245 }
2246 print_indentation(tabs, fp);
2247
2248 if (nir_loop_has_continue_construct(loop)) {
2249 fprintf(fp, "} continue {\n");
2250 foreach_list_typed(nir_cf_node, node, node, &loop->continue_list) {
2251 print_cf_node(node, state, tabs + 1);
2252 }
2253 print_indentation(tabs, fp);
2254 }
2255
2256 fprintf(fp, "}\n");
2257 }
2258
2259 static void
print_cf_node(nir_cf_node * node,print_state * state,unsigned int tabs)2260 print_cf_node(nir_cf_node *node, print_state *state, unsigned int tabs)
2261 {
2262 switch (node->type) {
2263 case nir_cf_node_block:
2264 print_block(nir_cf_node_as_block(node), state, tabs);
2265 break;
2266
2267 case nir_cf_node_if:
2268 print_if(nir_cf_node_as_if(node), state, tabs);
2269 break;
2270
2271 case nir_cf_node_loop:
2272 print_loop(nir_cf_node_as_loop(node), state, tabs);
2273 break;
2274
2275 default:
2276 unreachable("Invalid CFG node type");
2277 }
2278 }
2279
2280 static void
print_function_impl(nir_function_impl * impl,print_state * state)2281 print_function_impl(nir_function_impl *impl, print_state *state)
2282 {
2283 FILE *fp = state->fp;
2284
2285 state->max_dest_index = impl->ssa_alloc;
2286
2287 fprintf(fp, "\nimpl %s ", impl->function->name);
2288
2289 fprintf(fp, "{\n");
2290
2291 if (impl->preamble) {
2292 print_indentation(1, fp);
2293 fprintf(fp, "preamble %s\n", impl->preamble->name);
2294 }
2295
2296 if (!NIR_DEBUG(PRINT_NO_INLINE_CONSTS)) {
2297 /* Don't reindex the SSA as suggested by nir_gather_types() because
2298 * nir_print don't modify the shader. If needed, a limit for ssa_alloc
2299 * can be added.
2300 */
2301 state->float_types = calloc(BITSET_WORDS(impl->ssa_alloc), sizeof(BITSET_WORD));
2302 state->int_types = calloc(BITSET_WORDS(impl->ssa_alloc), sizeof(BITSET_WORD));
2303 nir_gather_types(impl, state->float_types, state->int_types);
2304 }
2305
2306 nir_foreach_function_temp_variable(var, impl) {
2307 print_indentation(1, fp);
2308 print_var_decl(var, state);
2309 }
2310
2311 nir_index_blocks(impl);
2312
2313 foreach_list_typed(nir_cf_node, node, node, &impl->body) {
2314 print_cf_node(node, state, 1);
2315 }
2316
2317 print_indentation(1, fp);
2318 fprintf(fp, "block b%u:\n}\n\n", impl->end_block->index);
2319
2320 free(state->float_types);
2321 free(state->int_types);
2322 state->max_dest_index = 0;
2323 }
2324
2325 static void
print_function(nir_function * function,print_state * state)2326 print_function(nir_function *function, print_state *state)
2327 {
2328 FILE *fp = state->fp;
2329
2330 fprintf(fp, "decl_function %s (", function->name);
2331
2332 for (unsigned i = 0; i < function->num_params; ++i) {
2333 if (i != 0) {
2334 fprintf(fp, ", ");
2335 }
2336
2337 nir_parameter param = function->params[i];
2338
2339 fprintf(fp, "%u", param.bit_size);
2340 if (param.num_components != 1) {
2341 fprintf(fp, "x%u", param.num_components);
2342 }
2343
2344 if (param.name) {
2345 fprintf(fp, " %s", param.name);
2346 } else if (param.is_return) {
2347 fprintf(fp, " return");
2348 }
2349 }
2350
2351 fprintf(fp, ")");
2352
2353 /* clang-format off */
2354 fprintf(fp, "%s%s%s", function->dont_inline ? " (noinline)" :
2355 function->should_inline ? " (inline)" : "",
2356 function->is_exported ? " (exported)" : "",
2357 function->is_entrypoint ? " (entrypoint)" : "");
2358 /* clang-format on */
2359
2360 if (function->workgroup_size[0]) {
2361 fprintf(fp, " (%ux%ux%u)",
2362 function->workgroup_size[0],
2363 function->workgroup_size[1],
2364 function->workgroup_size[2]);
2365 }
2366
2367 fprintf(fp, "\n");
2368
2369 if (function->impl != NULL) {
2370 print_function_impl(function->impl, state);
2371 return;
2372 }
2373 }
2374
2375 static void
init_print_state(print_state * state,nir_shader * shader,FILE * fp)2376 init_print_state(print_state *state, nir_shader *shader, FILE *fp)
2377 {
2378 state->fp = fp;
2379 state->shader = shader;
2380 state->ht = _mesa_pointer_hash_table_create(NULL);
2381 state->syms = _mesa_set_create(NULL, _mesa_hash_string,
2382 _mesa_key_string_equal);
2383 state->index = 0;
2384 state->int_types = NULL;
2385 state->float_types = NULL;
2386 state->max_dest_index = 0;
2387 state->padding_for_no_dest = 0;
2388 }
2389
2390 static void
destroy_print_state(print_state * state)2391 destroy_print_state(print_state *state)
2392 {
2393 _mesa_hash_table_destroy(state->ht, NULL);
2394 _mesa_set_destroy(state->syms, NULL);
2395 }
2396
2397 static const char *
primitive_name(unsigned primitive)2398 primitive_name(unsigned primitive)
2399 {
2400 #define PRIM(X) \
2401 case MESA_PRIM_##X: \
2402 return #X
2403 switch (primitive) {
2404 PRIM(POINTS);
2405 PRIM(LINES);
2406 PRIM(LINE_LOOP);
2407 PRIM(LINE_STRIP);
2408 PRIM(TRIANGLES);
2409 PRIM(TRIANGLE_STRIP);
2410 PRIM(TRIANGLE_FAN);
2411 PRIM(QUADS);
2412 PRIM(QUAD_STRIP);
2413 PRIM(POLYGON);
2414 PRIM(LINES_ADJACENCY);
2415 PRIM(TRIANGLES_ADJACENCY);
2416 default:
2417 return "UNKNOWN";
2418 }
2419 }
2420
2421 static void
print_bitset(FILE * fp,const char * label,const unsigned * words,int size)2422 print_bitset(FILE *fp, const char *label, const unsigned *words, int size)
2423 {
2424 fprintf(fp, "%s: ", label);
2425 /* Iterate back-to-front to get proper digit order (most significant first). */
2426 for (int i = size - 1; i >= 0; --i) {
2427 fprintf(fp, (i == size - 1) ? "0x%08x" : "'%08x", words[i]);
2428 }
2429 fprintf(fp, "\n");
2430 }
2431
2432 /* Print bitset, only if some bits are set */
2433 static void
print_nz_bitset(FILE * fp,const char * label,const unsigned * words,int size)2434 print_nz_bitset(FILE *fp, const char *label, const unsigned *words, int size)
2435 {
2436 bool is_all_zero = true;
2437 for (int i = 0; i < size; ++i) {
2438 if (words[i]) {
2439 is_all_zero = false;
2440 break;
2441 }
2442 }
2443
2444 if (!is_all_zero)
2445 print_bitset(fp, label, words, size);
2446 }
2447
2448 /* Print uint64_t value, only if non-zero.
2449 * The value is printed by enumerating the ranges of bits that are set.
2450 * E.g. inputs_read: 0,15-17
2451 */
2452 static void
print_nz_x64(FILE * fp,const char * label,uint64_t value)2453 print_nz_x64(FILE *fp, const char *label, uint64_t value)
2454 {
2455 if (value) {
2456 char acc[256] = { 0 };
2457 char buf[32];
2458 int start = 0;
2459 int count = 0;
2460 while (value) {
2461 u_bit_scan_consecutive_range64(&value, &start, &count);
2462 assert(count > 0);
2463 bool is_first = !acc[0];
2464 if (count > 1) {
2465 snprintf(buf, sizeof(buf), is_first ? "%d-%d" : ",%d-%d", start, start + count - 1);
2466 } else {
2467 snprintf(buf, sizeof(buf), is_first ? "%d" : ",%d", start);
2468 }
2469 assert(strlen(acc) + strlen(buf) + 1 < sizeof(acc));
2470 strcat(acc, buf);
2471 }
2472 fprintf(fp, "%s: %s\n", label, acc);
2473 }
2474 }
2475
2476 /* Print uint32_t value in hex, only if non-zero */
2477 static void
print_nz_x32(FILE * fp,const char * label,uint32_t value)2478 print_nz_x32(FILE *fp, const char *label, uint32_t value)
2479 {
2480 if (value)
2481 fprintf(fp, "%s: 0x%08" PRIx32 "\n", label, value);
2482 }
2483
2484 /* Print uint16_t value in hex, only if non-zero */
2485 static void
print_nz_x16(FILE * fp,const char * label,uint16_t value)2486 print_nz_x16(FILE *fp, const char *label, uint16_t value)
2487 {
2488 if (value)
2489 fprintf(fp, "%s: 0x%04x\n", label, value);
2490 }
2491
2492 /* Print uint8_t value in hex, only if non-zero */
2493 static void
print_nz_x8(FILE * fp,const char * label,uint8_t value)2494 print_nz_x8(FILE *fp, const char *label, uint8_t value)
2495 {
2496 if (value)
2497 fprintf(fp, "%s: 0x%02x\n", label, value);
2498 }
2499
2500 /* Print unsigned value in decimal, only if non-zero */
2501 static void
print_nz_unsigned(FILE * fp,const char * label,unsigned value)2502 print_nz_unsigned(FILE *fp, const char *label, unsigned value)
2503 {
2504 if (value)
2505 fprintf(fp, "%s: %u\n", label, value);
2506 }
2507
2508 /* Print bool only if set */
2509 static void
print_nz_bool(FILE * fp,const char * label,bool value)2510 print_nz_bool(FILE *fp, const char *label, bool value)
2511 {
2512 if (value)
2513 fprintf(fp, "%s: true\n", label);
2514 }
2515
2516 static void
print_shader_info(const struct shader_info * info,FILE * fp)2517 print_shader_info(const struct shader_info *info, FILE *fp)
2518 {
2519 fprintf(fp, "shader: %s\n", gl_shader_stage_name(info->stage));
2520
2521 if (memcmp(info->source_blake3, &(blake3_hash){0}, sizeof(info->source_blake3))) {
2522 fprintf(fp, "source_blake3: {");
2523 _mesa_blake3_print(fp, info->source_blake3);
2524 fprintf(fp, "}\n");
2525 }
2526
2527 if (info->name)
2528 fprintf(fp, "name: %s\n", info->name);
2529
2530 if (info->label)
2531 fprintf(fp, "label: %s\n", info->label);
2532
2533 print_nz_bool(fp, "internal", info->internal);
2534
2535 if (gl_shader_stage_uses_workgroup(info->stage)) {
2536 fprintf(fp, "workgroup_size: %u, %u, %u%s\n",
2537 info->workgroup_size[0],
2538 info->workgroup_size[1],
2539 info->workgroup_size[2],
2540 info->workgroup_size_variable ? " (variable)" : "");
2541 }
2542
2543 if (info->next_stage != MESA_SHADER_NONE)
2544 fprintf(fp, "next_stage: %s\n", gl_shader_stage_name(info->next_stage));
2545
2546 print_nz_unsigned(fp, "num_textures", info->num_textures);
2547 print_nz_unsigned(fp, "num_ubos", info->num_ubos);
2548 print_nz_unsigned(fp, "num_abos", info->num_abos);
2549 print_nz_unsigned(fp, "num_ssbos", info->num_ssbos);
2550 print_nz_unsigned(fp, "num_images", info->num_images);
2551
2552 print_nz_x64(fp, "inputs_read", info->inputs_read);
2553 print_nz_x64(fp, "dual_slot_inputs", info->dual_slot_inputs);
2554 print_nz_x64(fp, "outputs_written", info->outputs_written);
2555 print_nz_x64(fp, "outputs_read", info->outputs_read);
2556
2557 print_nz_bitset(fp, "system_values_read", info->system_values_read, ARRAY_SIZE(info->system_values_read));
2558
2559 print_nz_x64(fp, "per_primitive_inputs", info->per_primitive_inputs);
2560 print_nz_x64(fp, "per_primitive_outputs", info->per_primitive_outputs);
2561 print_nz_x64(fp, "per_view_outputs", info->per_view_outputs);
2562
2563 print_nz_x16(fp, "inputs_read_16bit", info->inputs_read_16bit);
2564 print_nz_x16(fp, "outputs_written_16bit", info->outputs_written_16bit);
2565 print_nz_x16(fp, "outputs_read_16bit", info->outputs_read_16bit);
2566 print_nz_x16(fp, "inputs_read_indirectly_16bit", info->inputs_read_indirectly_16bit);
2567 print_nz_x16(fp, "outputs_accessed_indirectly_16bit", info->outputs_accessed_indirectly_16bit);
2568
2569 print_nz_x32(fp, "patch_inputs_read", info->patch_inputs_read);
2570 print_nz_x32(fp, "patch_outputs_written", info->patch_outputs_written);
2571 print_nz_x32(fp, "patch_outputs_read", info->patch_outputs_read);
2572
2573 print_nz_x64(fp, "inputs_read_indirectly", info->inputs_read_indirectly);
2574 print_nz_x64(fp, "outputs_accessed_indirectly", info->outputs_accessed_indirectly);
2575 print_nz_x64(fp, "patch_inputs_read_indirectly", info->patch_inputs_read_indirectly);
2576 print_nz_x64(fp, "patch_outputs_accessed_indirectly", info->patch_outputs_accessed_indirectly);
2577
2578 print_nz_bitset(fp, "textures_used", info->textures_used, ARRAY_SIZE(info->textures_used));
2579 print_nz_bitset(fp, "textures_used_by_txf", info->textures_used_by_txf, ARRAY_SIZE(info->textures_used_by_txf));
2580 print_nz_bitset(fp, "samplers_used", info->samplers_used, ARRAY_SIZE(info->samplers_used));
2581 print_nz_bitset(fp, "images_used", info->images_used, ARRAY_SIZE(info->images_used));
2582 print_nz_bitset(fp, "image_buffers", info->image_buffers, ARRAY_SIZE(info->image_buffers));
2583 print_nz_bitset(fp, "msaa_images", info->msaa_images, ARRAY_SIZE(info->msaa_images));
2584
2585 print_nz_x32(fp, "float_controls_execution_mode", info->float_controls_execution_mode);
2586
2587 print_nz_unsigned(fp, "shared_size", info->shared_size);
2588
2589 if (info->stage == MESA_SHADER_MESH || info->stage == MESA_SHADER_TASK) {
2590 fprintf(fp, "task_payload_size: %u\n", info->task_payload_size);
2591 }
2592
2593 print_nz_unsigned(fp, "ray queries", info->ray_queries);
2594
2595 fprintf(fp, "subgroup_size: %u\n", info->subgroup_size);
2596
2597 print_nz_bool(fp, "uses_wide_subgroup_intrinsics", info->uses_wide_subgroup_intrinsics);
2598
2599 bool has_xfb_stride = info->xfb_stride[0] || info->xfb_stride[1] || info->xfb_stride[2] || info->xfb_stride[3];
2600 if (has_xfb_stride)
2601 fprintf(fp, "xfb_stride: {%u, %u, %u, %u}\n",
2602 info->xfb_stride[0],
2603 info->xfb_stride[1],
2604 info->xfb_stride[2],
2605 info->xfb_stride[3]);
2606
2607 bool has_inlinable_uniform_dw_offsets = info->inlinable_uniform_dw_offsets[0] || info->inlinable_uniform_dw_offsets[1] || info->inlinable_uniform_dw_offsets[2] || info->inlinable_uniform_dw_offsets[3];
2608 if (has_inlinable_uniform_dw_offsets)
2609 fprintf(fp, "inlinable_uniform_dw_offsets: {%u, %u, %u, %u}\n",
2610 info->inlinable_uniform_dw_offsets[0],
2611 info->inlinable_uniform_dw_offsets[1],
2612 info->inlinable_uniform_dw_offsets[2],
2613 info->inlinable_uniform_dw_offsets[3]);
2614
2615 print_nz_unsigned(fp, "num_inlinable_uniforms", info->num_inlinable_uniforms);
2616 print_nz_unsigned(fp, "clip_distance_array_size", info->clip_distance_array_size);
2617 print_nz_unsigned(fp, "cull_distance_array_size", info->cull_distance_array_size);
2618
2619 print_nz_bool(fp, "uses_texture_gather", info->uses_texture_gather);
2620 print_nz_bool(fp, "uses_resource_info_query", info->uses_resource_info_query);
2621 print_nz_bool(fp, "divergence_analysis_run", info->divergence_analysis_run);
2622
2623 print_nz_x8(fp, "bit_sizes_float", info->bit_sizes_float);
2624 print_nz_x8(fp, "bit_sizes_int", info->bit_sizes_int);
2625
2626 print_nz_bool(fp, "first_ubo_is_default_ubo", info->first_ubo_is_default_ubo);
2627 print_nz_bool(fp, "separate_shader", info->separate_shader);
2628 print_nz_bool(fp, "has_transform_feedback_varyings", info->has_transform_feedback_varyings);
2629 print_nz_bool(fp, "flrp_lowered", info->flrp_lowered);
2630 print_nz_bool(fp, "io_lowered", info->io_lowered);
2631 print_nz_bool(fp, "writes_memory", info->writes_memory);
2632 print_nz_unsigned(fp, "derivative_group", info->derivative_group);
2633
2634 switch (info->stage) {
2635 case MESA_SHADER_VERTEX:
2636 print_nz_x64(fp, "double_inputs", info->vs.double_inputs);
2637 print_nz_unsigned(fp, "blit_sgprs_amd", info->vs.blit_sgprs_amd);
2638 print_nz_bool(fp, "window_space_position", info->vs.window_space_position);
2639 print_nz_bool(fp, "needs_edge_flag", info->vs.needs_edge_flag);
2640 break;
2641
2642 case MESA_SHADER_TESS_CTRL:
2643 case MESA_SHADER_TESS_EVAL:
2644 fprintf(fp, "primitive_mode: %u\n", info->tess._primitive_mode);
2645 fprintf(fp, "tcs_vertices_out: %u\n", info->tess.tcs_vertices_out);
2646 fprintf(fp, "spacing: %u\n", info->tess.spacing);
2647
2648 print_nz_bool(fp, "ccw", info->tess.ccw);
2649 print_nz_bool(fp, "point_mode", info->tess.point_mode);
2650 print_nz_x64(fp, "tcs_same_invocation_inputs_read",
2651 info->tess.tcs_same_invocation_inputs_read);
2652 print_nz_x64(fp, "tcs_cross_invocation_inputs_read", info->tess.tcs_cross_invocation_inputs_read);
2653 print_nz_x64(fp, "tcs_cross_invocation_outputs_read", info->tess.tcs_cross_invocation_outputs_read);
2654 break;
2655
2656 case MESA_SHADER_GEOMETRY:
2657 fprintf(fp, "output_primitive: %s\n", primitive_name(info->gs.output_primitive));
2658 fprintf(fp, "input_primitive: %s\n", primitive_name(info->gs.input_primitive));
2659 fprintf(fp, "vertices_out: %u\n", info->gs.vertices_out);
2660 fprintf(fp, "invocations: %u\n", info->gs.invocations);
2661 fprintf(fp, "vertices_in: %u\n", info->gs.vertices_in);
2662 print_nz_bool(fp, "uses_end_primitive", info->gs.uses_end_primitive);
2663 fprintf(fp, "active_stream_mask: 0x%02x\n", info->gs.active_stream_mask);
2664 break;
2665
2666 case MESA_SHADER_FRAGMENT:
2667 print_nz_bool(fp, "uses_discard", info->fs.uses_discard);
2668 print_nz_bool(fp, "uses_fbfetch_output", info->fs.uses_fbfetch_output);
2669 print_nz_bool(fp, "color_is_dual_source", info->fs.color_is_dual_source);
2670
2671 print_nz_bool(fp, "require_full_quads", info->fs.require_full_quads);
2672 print_nz_bool(fp, "needs_quad_helper_invocations", info->fs.needs_quad_helper_invocations);
2673 print_nz_bool(fp, "uses_sample_qualifier", info->fs.uses_sample_qualifier);
2674 print_nz_bool(fp, "uses_sample_shading", info->fs.uses_sample_shading);
2675 print_nz_bool(fp, "early_fragment_tests", info->fs.early_fragment_tests);
2676 print_nz_bool(fp, "inner_coverage", info->fs.inner_coverage);
2677 print_nz_bool(fp, "post_depth_coverage", info->fs.post_depth_coverage);
2678
2679 print_nz_bool(fp, "pixel_center_integer", info->fs.pixel_center_integer);
2680 print_nz_bool(fp, "origin_upper_left", info->fs.origin_upper_left);
2681 print_nz_bool(fp, "pixel_interlock_ordered", info->fs.pixel_interlock_ordered);
2682 print_nz_bool(fp, "pixel_interlock_unordered", info->fs.pixel_interlock_unordered);
2683 print_nz_bool(fp, "sample_interlock_ordered", info->fs.sample_interlock_ordered);
2684 print_nz_bool(fp, "sample_interlock_unordered", info->fs.sample_interlock_unordered);
2685 print_nz_bool(fp, "untyped_color_outputs", info->fs.untyped_color_outputs);
2686
2687 print_nz_unsigned(fp, "depth_layout", info->fs.depth_layout);
2688
2689 if (info->fs.color0_interp != INTERP_MODE_NONE) {
2690 fprintf(fp, "color0_interp: %s\n",
2691 glsl_interp_mode_name(info->fs.color0_interp));
2692 }
2693 print_nz_bool(fp, "color0_sample", info->fs.color0_sample);
2694 print_nz_bool(fp, "color0_centroid", info->fs.color0_centroid);
2695
2696 if (info->fs.color1_interp != INTERP_MODE_NONE) {
2697 fprintf(fp, "color1_interp: %s\n",
2698 glsl_interp_mode_name(info->fs.color1_interp));
2699 }
2700 print_nz_bool(fp, "color1_sample", info->fs.color1_sample);
2701 print_nz_bool(fp, "color1_centroid", info->fs.color1_centroid);
2702
2703 print_nz_x32(fp, "advanced_blend_modes", info->fs.advanced_blend_modes);
2704 break;
2705
2706 case MESA_SHADER_COMPUTE:
2707 case MESA_SHADER_KERNEL:
2708 if (info->cs.workgroup_size_hint[0] || info->cs.workgroup_size_hint[1] || info->cs.workgroup_size_hint[2])
2709 fprintf(fp, "workgroup_size_hint: {%u, %u, %u}\n",
2710 info->cs.workgroup_size_hint[0],
2711 info->cs.workgroup_size_hint[1],
2712 info->cs.workgroup_size_hint[2]);
2713 print_nz_unsigned(fp, "user_data_components_amd", info->cs.user_data_components_amd);
2714 print_nz_unsigned(fp, "ptr_size", info->cs.ptr_size);
2715 break;
2716
2717 case MESA_SHADER_MESH:
2718 print_nz_x64(fp, "ms_cross_invocation_output_access", info->mesh.ms_cross_invocation_output_access);
2719 fprintf(fp, "max_vertices_out: %u\n", info->mesh.max_vertices_out);
2720 fprintf(fp, "max_primitives_out: %u\n", info->mesh.max_primitives_out);
2721 fprintf(fp, "primitive_type: %s\n", primitive_name(info->mesh.primitive_type));
2722 print_nz_bool(fp, "nv", info->mesh.nv);
2723 break;
2724
2725 default:
2726 fprintf(fp, "Unhandled stage %d\n", info->stage);
2727 }
2728 }
2729
2730 static void
_nir_print_shader_annotated(nir_shader * shader,FILE * fp,struct hash_table * annotations,nir_debug_info_instr ** debug_info)2731 _nir_print_shader_annotated(nir_shader *shader, FILE *fp,
2732 struct hash_table *annotations,
2733 nir_debug_info_instr **debug_info)
2734 {
2735 print_state state;
2736 init_print_state(&state, shader, fp);
2737 state.def_prefix = debug_info ? "ssa_" : "%";
2738 state.annotations = annotations;
2739 state.debug_info = debug_info;
2740
2741 print_shader_info(&shader->info, fp);
2742
2743 print_nz_unsigned(fp, "inputs", shader->num_inputs);
2744 print_nz_unsigned(fp, "outputs", shader->num_outputs);
2745 print_nz_unsigned(fp, "uniforms", shader->num_uniforms);
2746 if (shader->scratch_size)
2747 fprintf(fp, "scratch: %u\n", shader->scratch_size);
2748 if (shader->constant_data_size)
2749 fprintf(fp, "constants: %u\n", shader->constant_data_size);
2750 for (unsigned i = 0; i < nir_num_variable_modes; i++) {
2751 nir_variable_mode mode = BITFIELD_BIT(i);
2752 if (mode == nir_var_function_temp)
2753 continue;
2754
2755 if (mode == nir_var_shader_in || mode == nir_var_shader_out) {
2756 for (unsigned j = 0; j < 128; j++) {
2757 nir_variable *vars[NIR_MAX_VEC_COMPONENTS] = {0};
2758 nir_foreach_variable_with_modes(var, shader, mode) {
2759 if (var->data.location == j)
2760 vars[var->data.location_frac] = var;
2761 }
2762 for (unsigned j = 0; j < ARRAY_SIZE(vars); j++)
2763 if (vars[j]) {
2764 print_var_decl(vars[j], &state);
2765 }
2766 }
2767 } else {
2768 nir_foreach_variable_with_modes(var, shader, mode)
2769 print_var_decl(var, &state);
2770 }
2771 }
2772
2773 foreach_list_typed(nir_function, func, node, &shader->functions) {
2774 print_function(func, &state);
2775 }
2776
2777 destroy_print_state(&state);
2778 }
2779
2780 void
nir_print_shader_annotated(nir_shader * shader,FILE * fp,struct hash_table * annotations)2781 nir_print_shader_annotated(nir_shader *shader, FILE *fp,
2782 struct hash_table *annotations)
2783 {
2784 _nir_print_shader_annotated(shader, fp, annotations, NULL);
2785 }
2786
2787 void
nir_print_shader(nir_shader * shader,FILE * fp)2788 nir_print_shader(nir_shader *shader, FILE *fp)
2789 {
2790 nir_print_shader_annotated(shader, fp, NULL);
2791 fflush(fp);
2792 }
2793
2794 static char *
_nir_shader_as_str_annotated(nir_shader * nir,struct hash_table * annotations,void * mem_ctx,nir_debug_info_instr ** debug_info)2795 _nir_shader_as_str_annotated(nir_shader *nir, struct hash_table *annotations, void *mem_ctx,
2796 nir_debug_info_instr **debug_info)
2797 {
2798 char *stream_data = NULL;
2799 size_t stream_size = 0;
2800 struct u_memstream mem;
2801 if (u_memstream_open(&mem, &stream_data, &stream_size)) {
2802 FILE *const stream = u_memstream_get(&mem);
2803 _nir_print_shader_annotated(nir, stream, annotations, debug_info);
2804 u_memstream_close(&mem);
2805 }
2806
2807 char *str = ralloc_size(mem_ctx, stream_size + 1);
2808 memcpy(str, stream_data, stream_size);
2809 str[stream_size] = '\0';
2810
2811 free(stream_data);
2812
2813 return str;
2814 }
2815
2816 char *
nir_shader_as_str_annotated(nir_shader * nir,struct hash_table * annotations,void * mem_ctx)2817 nir_shader_as_str_annotated(nir_shader *nir, struct hash_table *annotations, void *mem_ctx)
2818 {
2819 return _nir_shader_as_str_annotated(nir, annotations, mem_ctx, NULL);
2820 }
2821
2822 char *
nir_shader_as_str(nir_shader * nir,void * mem_ctx)2823 nir_shader_as_str(nir_shader *nir, void *mem_ctx)
2824 {
2825 return nir_shader_as_str_annotated(nir, NULL, mem_ctx);
2826 }
2827
2828 void
nir_print_instr(const nir_instr * instr,FILE * fp)2829 nir_print_instr(const nir_instr *instr, FILE *fp)
2830 {
2831 print_state state = {
2832 .fp = fp,
2833 .def_prefix = "%",
2834 };
2835 if (instr->block) {
2836 nir_function_impl *impl = nir_cf_node_get_function(&instr->block->cf_node);
2837 state.shader = impl->function->shader;
2838 }
2839
2840 print_instr(instr, &state, 0);
2841 }
2842
2843 char *
nir_instr_as_str(const nir_instr * instr,void * mem_ctx)2844 nir_instr_as_str(const nir_instr *instr, void *mem_ctx)
2845 {
2846 char *stream_data = NULL;
2847 size_t stream_size = 0;
2848 struct u_memstream mem;
2849 if (u_memstream_open(&mem, &stream_data, &stream_size)) {
2850 FILE *const stream = u_memstream_get(&mem);
2851 nir_print_instr(instr, stream);
2852 u_memstream_close(&mem);
2853 }
2854
2855 char *str = ralloc_size(mem_ctx, stream_size + 1);
2856 memcpy(str, stream_data, stream_size);
2857 str[stream_size] = '\0';
2858
2859 free(stream_data);
2860
2861 return str;
2862 }
2863
2864 void
nir_print_deref(const nir_deref_instr * deref,FILE * fp)2865 nir_print_deref(const nir_deref_instr *deref, FILE *fp)
2866 {
2867 print_state state = {
2868 .fp = fp,
2869 .def_prefix = "%",
2870 };
2871 print_deref_link(deref, true, &state);
2872 }
2873
2874 void
nir_log_shader_annotated_tagged(enum mesa_log_level level,const char * tag,nir_shader * shader,struct hash_table * annotations)2875 nir_log_shader_annotated_tagged(enum mesa_log_level level, const char *tag,
2876 nir_shader *shader, struct hash_table *annotations)
2877 {
2878 char *str = nir_shader_as_str_annotated(shader, annotations, NULL);
2879 _mesa_log_multiline(level, tag, str);
2880 ralloc_free(str);
2881 }
2882
2883 char *
nir_shader_gather_debug_info(nir_shader * shader,const char * filename,uint32_t first_line)2884 nir_shader_gather_debug_info(nir_shader *shader, const char *filename, uint32_t first_line)
2885 {
2886 uint32_t instr_count = 0;
2887 nir_foreach_function_impl(impl, shader) {
2888 nir_foreach_block(block, impl) {
2889 nir_foreach_instr(instr, block) {
2890 instr->index = instr_count;
2891 instr_count++;
2892 }
2893 }
2894 }
2895
2896 if (!instr_count)
2897 return nir_shader_as_str(shader, NULL);
2898
2899 nir_debug_info_instr **debug_info = rzalloc_array(shader, nir_debug_info_instr *, instr_count);
2900
2901 instr_count = 0;
2902 nir_foreach_function_impl(impl, shader) {
2903 nir_builder b = nir_builder_at(nir_before_cf_list(&impl->body));
2904 nir_def *filename_def = nir_build_string(&b, filename);
2905
2906 nir_foreach_block(block, impl) {
2907 nir_foreach_instr_safe(instr, block) {
2908 if (instr->type == nir_instr_type_debug_info ||
2909 instr->type == nir_instr_type_phi)
2910 continue;
2911
2912 nir_debug_info_instr *di = nir_debug_info_instr_create(shader, nir_debug_info_src_loc, 0);
2913 di->src_loc.filename = nir_src_for_ssa(filename_def);
2914 di->src_loc.source = nir_debug_info_nir;
2915 debug_info[instr_count++] = di;
2916 }
2917 }
2918 }
2919
2920 char *str = _nir_shader_as_str_annotated(shader, NULL, NULL, debug_info);
2921
2922 uint32_t line = first_line;
2923 uint32_t character_index = 0;
2924
2925 for (uint32_t i = 0; i < instr_count; i++) {
2926 nir_debug_info_instr *di = debug_info[i];
2927 if (!di)
2928 continue;
2929
2930 while (character_index < di->src_loc.column) {
2931 if (str[character_index] == '\n')
2932 line++;
2933 character_index++;
2934 }
2935
2936 di->src_loc.line = line;
2937 di->src_loc.column = 0;
2938 }
2939
2940 instr_count = 0;
2941 nir_foreach_function_impl(impl, shader) {
2942 nir_foreach_block(block, impl) {
2943 nir_foreach_instr_safe(instr, block) {
2944 if (instr->type != nir_instr_type_debug_info &&
2945 instr->type != nir_instr_type_phi)
2946 nir_instr_insert_before(instr, &debug_info[instr_count++]->instr);
2947 }
2948 }
2949 }
2950
2951 return str;
2952 }
2953