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 ([email protected])
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 == 1) + 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_VOLATILE, "volatile" },
799 { ACCESS_RESTRICT, "restrict" },
800 { ACCESS_NON_WRITEABLE, "readonly" },
801 { ACCESS_NON_READABLE, "writeonly" },
802 { ACCESS_CAN_REORDER, "reorderable" },
803 { ACCESS_CAN_SPECULATE, "speculatable" },
804 { ACCESS_NON_TEMPORAL, "non-temporal" },
805 { ACCESS_INCLUDE_HELPERS, "include-helpers" },
806 { ACCESS_CP_GE_COHERENT_AMD, "cp-ge-coherent-amd" },
807 };
808
809 bool first = true;
810 for (unsigned i = 0; i < ARRAY_SIZE(modes); ++i) {
811 if (access & modes[i].bit) {
812 fprintf(state->fp, "%s%s", first ? "" : separator, modes[i].name);
813 first = false;
814 }
815 }
816 }
817
818 static void
print_var_decl(nir_variable * var,print_state * state)819 print_var_decl(nir_variable *var, print_state *state)
820 {
821 FILE *fp = state->fp;
822
823 fprintf(fp, "decl_var ");
824
825 const char *const bindless = (var->data.bindless) ? "bindless " : "";
826 const char *const cent = (var->data.centroid) ? "centroid " : "";
827 const char *const samp = (var->data.sample) ? "sample " : "";
828 const char *const patch = (var->data.patch) ? "patch " : "";
829 const char *const inv = (var->data.invariant) ? "invariant " : "";
830 const char *const per_view = (var->data.per_view) ? "per_view " : "";
831 const char *const per_primitive = (var->data.per_primitive) ? "per_primitive " : "";
832 const char *const ray_query = (var->data.ray_query) ? "ray_query " : "";
833 fprintf(fp, "%s%s%s%s%s%s%s%s%s %s ",
834 bindless, cent, samp, patch, inv, per_view, per_primitive, ray_query,
835 get_variable_mode_str(var->data.mode, false),
836 glsl_interp_mode_name(var->data.interpolation));
837
838 print_access(var->data.access, state, " ");
839 fprintf(fp, " ");
840
841 if (glsl_get_base_type(glsl_without_array(var->type)) == GLSL_TYPE_IMAGE) {
842 fprintf(fp, "%s ", util_format_short_name(var->data.image.format));
843 }
844
845 if (var->data.precision) {
846 const char *precisions[] = {
847 "",
848 "highp",
849 "mediump",
850 "lowp",
851 };
852 fprintf(fp, "%s ", precisions[var->data.precision]);
853 }
854
855 fprintf(fp, "%s %s", glsl_get_type_name(var->type),
856 get_var_name(var, state));
857
858 if (var->data.mode & (nir_var_shader_in |
859 nir_var_shader_out |
860 nir_var_uniform |
861 nir_var_system_value |
862 nir_var_mem_ubo |
863 nir_var_mem_ssbo |
864 nir_var_image)) {
865 char buf[4];
866 const char *loc = get_location_str(var->data.location,
867 state->shader->info.stage,
868 var->data.mode, buf);
869
870 /* For shader I/O vars that have been split to components or packed,
871 * print the fractional location within the input/output.
872 */
873 unsigned int num_components =
874 glsl_get_components(glsl_without_array(var->type));
875 const char *components = "";
876 char components_local[18] = { '.' /* the rest is 0-filled */ };
877 switch (var->data.mode) {
878 case nir_var_shader_in:
879 case nir_var_shader_out:
880 if (num_components < 16 && num_components != 0) {
881 const char *xyzw = comp_mask_string(num_components);
882 for (int i = 0; i < num_components; i++)
883 components_local[i + 1] = xyzw[i + var->data.location_frac];
884
885 components = components_local;
886 }
887 break;
888 default:
889 break;
890 }
891
892 if (var->data.mode & nir_var_system_value) {
893 fprintf(fp, " (%s%s)", loc, components);
894 } else {
895 fprintf(fp, " (%s%s, %u, %u)%s", loc,
896 components,
897 var->data.driver_location, var->data.binding,
898 var->data.compact ? " compact" : "");
899 }
900 }
901
902 if (var->constant_initializer) {
903 if (var->constant_initializer->is_null_constant) {
904 fprintf(fp, " = null");
905 } else {
906 fprintf(fp, " = { ");
907 print_constant(var->constant_initializer, var->type, state);
908 fprintf(fp, " }");
909 }
910 }
911 if (glsl_type_is_sampler(var->type) && var->data.sampler.is_inline_sampler) {
912 fprintf(fp, " = { %s, %s, %s }",
913 get_constant_sampler_addressing_mode(var->data.sampler.addressing_mode),
914 var->data.sampler.normalized_coordinates ? "true" : "false",
915 get_constant_sampler_filter_mode(var->data.sampler.filter_mode));
916 }
917 if (var->pointer_initializer)
918 fprintf(fp, " = &%s", get_var_name(var->pointer_initializer, state));
919
920 fprintf(fp, "\n");
921 print_annotation(state, var);
922 }
923
924 static void
print_deref_link(const nir_deref_instr * instr,bool whole_chain,print_state * state)925 print_deref_link(const nir_deref_instr *instr, bool whole_chain, print_state *state)
926 {
927 FILE *fp = state->fp;
928
929 if (instr->deref_type == nir_deref_type_var) {
930 fprintf(fp, "%s", get_var_name(instr->var, state));
931 return;
932 } else if (instr->deref_type == nir_deref_type_cast) {
933 fprintf(fp, "(%s *)", glsl_get_type_name(instr->type));
934 print_src(&instr->parent, state, nir_type_invalid);
935 return;
936 }
937
938 nir_deref_instr *parent =
939 nir_instr_as_deref(instr->parent.ssa->parent_instr);
940
941 /* Is the parent we're going to print a bare cast? */
942 const bool is_parent_cast =
943 whole_chain && parent->deref_type == nir_deref_type_cast;
944
945 /* If we're not printing the whole chain, the parent we print will be a SSA
946 * value that represents a pointer. The only deref type that naturally
947 * gives a pointer is a cast.
948 */
949 const bool is_parent_pointer =
950 !whole_chain || parent->deref_type == nir_deref_type_cast;
951
952 /* Struct derefs have a nice syntax that works on pointers, arrays derefs
953 * do not.
954 */
955 const bool need_deref =
956 is_parent_pointer && instr->deref_type != nir_deref_type_struct;
957
958 /* Cast need extra parens and so * dereferences */
959 if (is_parent_cast || need_deref)
960 fprintf(fp, "(");
961
962 if (need_deref)
963 fprintf(fp, "*");
964
965 if (whole_chain) {
966 print_deref_link(parent, whole_chain, state);
967 } else {
968 print_src(&instr->parent, state, nir_type_invalid);
969 }
970
971 if (is_parent_cast || need_deref)
972 fprintf(fp, ")");
973
974 switch (instr->deref_type) {
975 case nir_deref_type_struct:
976 fprintf(fp, "%s%s", is_parent_pointer ? "->" : ".",
977 glsl_get_struct_elem_name(parent->type, instr->strct.index));
978 break;
979
980 case nir_deref_type_array:
981 case nir_deref_type_ptr_as_array: {
982 if (nir_src_is_const(instr->arr.index)) {
983 fprintf(fp, "[%" PRId64 "]", nir_src_as_int(instr->arr.index));
984 } else {
985 fprintf(fp, "[");
986 print_src(&instr->arr.index, state, nir_type_invalid);
987 fprintf(fp, "]");
988 }
989 break;
990 }
991
992 case nir_deref_type_array_wildcard:
993 fprintf(fp, "[*]");
994 break;
995
996 default:
997 unreachable("Invalid deref instruction type");
998 }
999 }
1000
1001 static void
print_deref_instr(nir_deref_instr * instr,print_state * state)1002 print_deref_instr(nir_deref_instr *instr, print_state *state)
1003 {
1004 FILE *fp = state->fp;
1005
1006 print_def(&instr->def, state);
1007
1008 switch (instr->deref_type) {
1009 case nir_deref_type_var:
1010 fprintf(fp, " = deref_var ");
1011 break;
1012 case nir_deref_type_array:
1013 case nir_deref_type_array_wildcard:
1014 fprintf(fp, " = deref_array ");
1015 break;
1016 case nir_deref_type_struct:
1017 fprintf(fp, " = deref_struct ");
1018 break;
1019 case nir_deref_type_cast:
1020 fprintf(fp, " = deref_cast ");
1021 break;
1022 case nir_deref_type_ptr_as_array:
1023 fprintf(fp, " = deref_ptr_as_array ");
1024 break;
1025 default:
1026 unreachable("Invalid deref instruction type");
1027 }
1028
1029 /* Only casts naturally return a pointer type */
1030 if (instr->deref_type != nir_deref_type_cast)
1031 fprintf(fp, "&");
1032
1033 print_deref_link(instr, false, state);
1034
1035 fprintf(fp, " (");
1036 unsigned modes = instr->modes;
1037 while (modes) {
1038 int m = u_bit_scan(&modes);
1039 fprintf(fp, "%s%s", get_variable_mode_str(1 << m, true),
1040 modes ? "|" : "");
1041 }
1042 fprintf(fp, " %s)", glsl_get_type_name(instr->type));
1043
1044 if (instr->deref_type == nir_deref_type_cast) {
1045 fprintf(fp, " (ptr_stride=%u, align_mul=%u, align_offset=%u)",
1046 instr->cast.ptr_stride,
1047 instr->cast.align_mul, instr->cast.align_offset);
1048 }
1049
1050 if (instr->deref_type != nir_deref_type_var &&
1051 instr->deref_type != nir_deref_type_cast) {
1052 /* Print the entire chain as a comment */
1053 fprintf(fp, " // &");
1054 print_deref_link(instr, true, state);
1055 }
1056 }
1057
1058 static const char *
vulkan_descriptor_type_name(VkDescriptorType type)1059 vulkan_descriptor_type_name(VkDescriptorType type)
1060 {
1061 switch (type) {
1062 case VK_DESCRIPTOR_TYPE_SAMPLER:
1063 return "sampler";
1064 case VK_DESCRIPTOR_TYPE_COMBINED_IMAGE_SAMPLER:
1065 return "texture+sampler";
1066 case VK_DESCRIPTOR_TYPE_SAMPLED_IMAGE:
1067 return "texture";
1068 case VK_DESCRIPTOR_TYPE_STORAGE_IMAGE:
1069 return "image";
1070 case VK_DESCRIPTOR_TYPE_UNIFORM_TEXEL_BUFFER:
1071 return "texture-buffer";
1072 case VK_DESCRIPTOR_TYPE_STORAGE_TEXEL_BUFFER:
1073 return "image-buffer";
1074 case VK_DESCRIPTOR_TYPE_UNIFORM_BUFFER:
1075 return "UBO";
1076 case VK_DESCRIPTOR_TYPE_STORAGE_BUFFER:
1077 return "SSBO";
1078 case VK_DESCRIPTOR_TYPE_UNIFORM_BUFFER_DYNAMIC:
1079 return "UBO";
1080 case VK_DESCRIPTOR_TYPE_STORAGE_BUFFER_DYNAMIC:
1081 return "SSBO";
1082 case VK_DESCRIPTOR_TYPE_INPUT_ATTACHMENT:
1083 return "input-att";
1084 case VK_DESCRIPTOR_TYPE_INLINE_UNIFORM_BLOCK:
1085 return "inline-UBO";
1086 case VK_DESCRIPTOR_TYPE_ACCELERATION_STRUCTURE_KHR:
1087 return "accel-struct";
1088 default:
1089 return "unknown";
1090 }
1091 }
1092
1093 static void
print_alu_type(nir_alu_type type,print_state * state)1094 print_alu_type(nir_alu_type type, print_state *state)
1095 {
1096 FILE *fp = state->fp;
1097 unsigned size = nir_alu_type_get_type_size(type);
1098 const char *name;
1099
1100 switch (nir_alu_type_get_base_type(type)) {
1101 case nir_type_int:
1102 name = "int";
1103 break;
1104 case nir_type_uint:
1105 name = "uint";
1106 break;
1107 case nir_type_bool:
1108 name = "bool";
1109 break;
1110 case nir_type_float:
1111 name = "float";
1112 break;
1113 default:
1114 name = "invalid";
1115 }
1116 if (size)
1117 fprintf(fp, "%s%u", name, size);
1118 else
1119 fprintf(fp, "%s", name);
1120 }
1121
1122 static void
print_intrinsic_instr(nir_intrinsic_instr * instr,print_state * state)1123 print_intrinsic_instr(nir_intrinsic_instr *instr, print_state *state)
1124 {
1125 const nir_intrinsic_info *info = &nir_intrinsic_infos[instr->intrinsic];
1126 unsigned num_srcs = info->num_srcs;
1127 FILE *fp = state->fp;
1128
1129 if (info->has_dest) {
1130 print_def(&instr->def, state);
1131 fprintf(fp, " = ");
1132 } else {
1133 print_no_dest_padding(state);
1134 }
1135
1136 fprintf(fp, "@%s", info->name);
1137
1138 for (unsigned i = 0; i < num_srcs; i++) {
1139 if (i == 0)
1140 fprintf(fp, " (");
1141 else
1142 fprintf(fp, ", ");
1143
1144 print_src(&instr->src[i], state, nir_intrinsic_instr_src_type(instr, i));
1145 }
1146
1147 if (num_srcs)
1148 fprintf(fp, ")");
1149
1150 for (unsigned i = 0; i < info->num_indices; i++) {
1151 unsigned idx = info->indices[i];
1152 if (i == 0)
1153 fprintf(fp, " (");
1154 else
1155 fprintf(fp, ", ");
1156 switch (idx) {
1157 case NIR_INTRINSIC_WRITE_MASK: {
1158 /* special case wrmask to show it as a writemask.. */
1159 unsigned wrmask = nir_intrinsic_write_mask(instr);
1160 fprintf(fp, "wrmask=");
1161 for (unsigned i = 0; i < instr->num_components; i++)
1162 if ((wrmask >> i) & 1)
1163 fprintf(fp, "%c", comp_mask_string(instr->num_components)[i]);
1164 break;
1165 }
1166
1167 case NIR_INTRINSIC_REDUCTION_OP: {
1168 nir_op reduction_op = nir_intrinsic_reduction_op(instr);
1169 fprintf(fp, "reduction_op=%s", nir_op_infos[reduction_op].name);
1170 break;
1171 }
1172
1173 case NIR_INTRINSIC_ATOMIC_OP: {
1174 nir_atomic_op atomic_op = nir_intrinsic_atomic_op(instr);
1175 fprintf(fp, "atomic_op=");
1176
1177 switch (atomic_op) {
1178 case nir_atomic_op_iadd:
1179 fprintf(fp, "iadd");
1180 break;
1181 case nir_atomic_op_imin:
1182 fprintf(fp, "imin");
1183 break;
1184 case nir_atomic_op_umin:
1185 fprintf(fp, "umin");
1186 break;
1187 case nir_atomic_op_imax:
1188 fprintf(fp, "imax");
1189 break;
1190 case nir_atomic_op_umax:
1191 fprintf(fp, "umax");
1192 break;
1193 case nir_atomic_op_iand:
1194 fprintf(fp, "iand");
1195 break;
1196 case nir_atomic_op_ior:
1197 fprintf(fp, "ior");
1198 break;
1199 case nir_atomic_op_ixor:
1200 fprintf(fp, "ixor");
1201 break;
1202 case nir_atomic_op_xchg:
1203 fprintf(fp, "xchg");
1204 break;
1205 case nir_atomic_op_fadd:
1206 fprintf(fp, "fadd");
1207 break;
1208 case nir_atomic_op_fmin:
1209 fprintf(fp, "fmin");
1210 break;
1211 case nir_atomic_op_fmax:
1212 fprintf(fp, "fmax");
1213 break;
1214 case nir_atomic_op_cmpxchg:
1215 fprintf(fp, "cmpxchg");
1216 break;
1217 case nir_atomic_op_fcmpxchg:
1218 fprintf(fp, "fcmpxchg");
1219 break;
1220 case nir_atomic_op_inc_wrap:
1221 fprintf(fp, "inc_wrap");
1222 break;
1223 case nir_atomic_op_dec_wrap:
1224 fprintf(fp, "dec_wrap");
1225 break;
1226 case nir_atomic_op_ordered_add_gfx12_amd:
1227 fprintf(fp, "ordered_add");
1228 break;
1229 }
1230 break;
1231 }
1232
1233 case NIR_INTRINSIC_IMAGE_DIM: {
1234 static const char *dim_name[] = {
1235 [GLSL_SAMPLER_DIM_1D] = "1D",
1236 [GLSL_SAMPLER_DIM_2D] = "2D",
1237 [GLSL_SAMPLER_DIM_3D] = "3D",
1238 [GLSL_SAMPLER_DIM_CUBE] = "Cube",
1239 [GLSL_SAMPLER_DIM_RECT] = "Rect",
1240 [GLSL_SAMPLER_DIM_BUF] = "Buf",
1241 [GLSL_SAMPLER_DIM_MS] = "2D-MSAA",
1242 [GLSL_SAMPLER_DIM_SUBPASS] = "Subpass",
1243 [GLSL_SAMPLER_DIM_SUBPASS_MS] = "Subpass-MSAA",
1244 };
1245 enum glsl_sampler_dim dim = nir_intrinsic_image_dim(instr);
1246 assert(dim < ARRAY_SIZE(dim_name) && dim_name[dim]);
1247 fprintf(fp, "image_dim=%s", dim_name[dim]);
1248 break;
1249 }
1250
1251 case NIR_INTRINSIC_IMAGE_ARRAY: {
1252 bool array = nir_intrinsic_image_array(instr);
1253 fprintf(fp, "image_array=%s", array ? "true" : "false");
1254 break;
1255 }
1256
1257 case NIR_INTRINSIC_FORMAT: {
1258 enum pipe_format format = nir_intrinsic_format(instr);
1259 fprintf(fp, "format=%s", util_format_short_name(format));
1260 break;
1261 }
1262
1263 case NIR_INTRINSIC_DESC_TYPE: {
1264 VkDescriptorType desc_type = nir_intrinsic_desc_type(instr);
1265 fprintf(fp, "desc_type=%s", vulkan_descriptor_type_name(desc_type));
1266 break;
1267 }
1268
1269 case NIR_INTRINSIC_SRC_TYPE: {
1270 fprintf(fp, "src_type=");
1271 print_alu_type(nir_intrinsic_src_type(instr), state);
1272 break;
1273 }
1274
1275 case NIR_INTRINSIC_DEST_TYPE: {
1276 fprintf(fp, "dest_type=");
1277 print_alu_type(nir_intrinsic_dest_type(instr), state);
1278 break;
1279 }
1280
1281 case NIR_INTRINSIC_SWIZZLE_MASK: {
1282 fprintf(fp, "swizzle_mask=");
1283 unsigned mask = nir_intrinsic_swizzle_mask(instr);
1284 if (instr->intrinsic == nir_intrinsic_quad_swizzle_amd) {
1285 for (unsigned i = 0; i < 4; i++)
1286 fprintf(fp, "%d", (mask >> (i * 2) & 3));
1287 } else if (instr->intrinsic == nir_intrinsic_masked_swizzle_amd) {
1288 fprintf(fp, "((id & %d) | %d) ^ %d", mask & 0x1F,
1289 (mask >> 5) & 0x1F,
1290 (mask >> 10) & 0x1F);
1291 } else {
1292 fprintf(fp, "%d", mask);
1293 }
1294 break;
1295 }
1296
1297 case NIR_INTRINSIC_MEMORY_SEMANTICS: {
1298 nir_memory_semantics semantics = nir_intrinsic_memory_semantics(instr);
1299 fprintf(fp, "mem_semantics=");
1300 switch (semantics & (NIR_MEMORY_ACQUIRE | NIR_MEMORY_RELEASE)) {
1301 case 0:
1302 fprintf(fp, "NONE");
1303 break;
1304 case NIR_MEMORY_ACQUIRE:
1305 fprintf(fp, "ACQ");
1306 break;
1307 case NIR_MEMORY_RELEASE:
1308 fprintf(fp, "REL");
1309 break;
1310 default:
1311 fprintf(fp, "ACQ|REL");
1312 break;
1313 }
1314 if (semantics & (NIR_MEMORY_MAKE_AVAILABLE))
1315 fprintf(fp, "|AVAILABLE");
1316 if (semantics & (NIR_MEMORY_MAKE_VISIBLE))
1317 fprintf(fp, "|VISIBLE");
1318 break;
1319 }
1320
1321 case NIR_INTRINSIC_MEMORY_MODES: {
1322 fprintf(fp, "mem_modes=");
1323 unsigned int modes = nir_intrinsic_memory_modes(instr);
1324 if (modes == 0)
1325 fputc('0', fp);
1326 while (modes) {
1327 nir_variable_mode m = u_bit_scan(&modes);
1328 fprintf(fp, "%s%s", get_variable_mode_str(1 << m, true), modes ? "|" : "");
1329 }
1330 break;
1331 }
1332
1333 case NIR_INTRINSIC_EXECUTION_SCOPE:
1334 case NIR_INTRINSIC_MEMORY_SCOPE: {
1335 mesa_scope scope =
1336 idx == NIR_INTRINSIC_MEMORY_SCOPE ? nir_intrinsic_memory_scope(instr)
1337 : nir_intrinsic_execution_scope(instr);
1338 const char *name = mesa_scope_name(scope);
1339 static const char prefix[] = "SCOPE_";
1340 if (strncmp(name, prefix, sizeof(prefix) - 1) == 0)
1341 name += sizeof(prefix) - 1;
1342 fprintf(fp, "%s=%s", nir_intrinsic_index_names[idx], name);
1343 break;
1344 }
1345
1346 case NIR_INTRINSIC_IO_SEMANTICS: {
1347 struct nir_io_semantics io = nir_intrinsic_io_semantics(instr);
1348
1349 /* Try to figure out the mode so we can interpret the location */
1350 nir_variable_mode mode = nir_var_mem_generic;
1351 switch (instr->intrinsic) {
1352 case nir_intrinsic_load_input:
1353 case nir_intrinsic_load_per_primitive_input:
1354 case nir_intrinsic_load_interpolated_input:
1355 case nir_intrinsic_load_per_vertex_input:
1356 case nir_intrinsic_load_input_vertex:
1357 case nir_intrinsic_load_coefficients_agx:
1358 mode = nir_var_shader_in;
1359 break;
1360
1361 case nir_intrinsic_load_output:
1362 case nir_intrinsic_store_output:
1363 case nir_intrinsic_store_per_primitive_output:
1364 case nir_intrinsic_store_per_vertex_output:
1365 mode = nir_var_shader_out;
1366 break;
1367
1368 default:
1369 break;
1370 }
1371
1372 /* Using that mode, we should be able to name the location */
1373 char buf[4];
1374 const char *loc = get_location_str(io.location,
1375 state->shader->info.stage, mode,
1376 buf);
1377
1378 fprintf(fp, "io location=%s slots=%u", loc, io.num_slots);
1379
1380 if (io.interp_explicit_strict)
1381 fprintf(fp, " explicit_strict");
1382
1383 if (io.dual_source_blend_index)
1384 fprintf(fp, " dualsrc");
1385
1386 if (io.fb_fetch_output)
1387 fprintf(fp, " fbfetch");
1388
1389 if (io.per_view)
1390 fprintf(fp, " perview");
1391
1392 if (io.medium_precision)
1393 fprintf(fp, " mediump");
1394
1395 if (io.high_16bits)
1396 fprintf(fp, " high_16bits");
1397
1398 if (io.invariant)
1399 fprintf(fp, " invariant");
1400
1401 if (io.high_dvec2)
1402 fprintf(fp, " high_dvec2");
1403
1404 if (io.no_varying)
1405 fprintf(fp, " no_varying");
1406
1407 if (io.no_sysval_output)
1408 fprintf(fp, " no_sysval_output");
1409
1410 if (state->shader &&
1411 state->shader->info.stage == MESA_SHADER_GEOMETRY &&
1412 (instr->intrinsic == nir_intrinsic_store_output ||
1413 instr->intrinsic == nir_intrinsic_store_per_primitive_output ||
1414 instr->intrinsic == nir_intrinsic_store_per_vertex_output)) {
1415 unsigned gs_streams = io.gs_streams;
1416 fprintf(fp, " gs_streams(");
1417 for (unsigned i = 0; i < 4; i++) {
1418 fprintf(fp, "%s%c=%u", i ? " " : "", "xyzw"[i],
1419 (gs_streams >> (i * 2)) & 0x3);
1420 }
1421 fprintf(fp, ")");
1422 }
1423
1424 break;
1425 }
1426
1427 case NIR_INTRINSIC_IO_XFB:
1428 case NIR_INTRINSIC_IO_XFB2: {
1429 /* This prints both IO_XFB and IO_XFB2. */
1430 fprintf(fp, "xfb%s(", idx == NIR_INTRINSIC_IO_XFB ? "" : "2");
1431 bool first = true;
1432 for (unsigned i = 0; i < 2; i++) {
1433 unsigned start_comp = (idx == NIR_INTRINSIC_IO_XFB ? 0 : 2) + i;
1434 nir_io_xfb xfb = start_comp < 2 ? nir_intrinsic_io_xfb(instr) : nir_intrinsic_io_xfb2(instr);
1435
1436 if (!xfb.out[i].num_components)
1437 continue;
1438
1439 if (!first)
1440 fprintf(fp, ", ");
1441 first = false;
1442
1443 if (xfb.out[i].num_components > 1) {
1444 fprintf(fp, "components=%u..%u",
1445 start_comp, start_comp + xfb.out[i].num_components - 1);
1446 } else {
1447 fprintf(fp, "component=%u", start_comp);
1448 }
1449 fprintf(fp, " buffer=%u offset=%u",
1450 xfb.out[i].buffer, (uint32_t)xfb.out[i].offset * 4);
1451 }
1452 fprintf(fp, ")");
1453 break;
1454 }
1455
1456 case NIR_INTRINSIC_ROUNDING_MODE: {
1457 fprintf(fp, "rounding_mode=");
1458 switch (nir_intrinsic_rounding_mode(instr)) {
1459 case nir_rounding_mode_undef:
1460 fprintf(fp, "undef");
1461 break;
1462 case nir_rounding_mode_rtne:
1463 fprintf(fp, "rtne");
1464 break;
1465 case nir_rounding_mode_ru:
1466 fprintf(fp, "ru");
1467 break;
1468 case nir_rounding_mode_rd:
1469 fprintf(fp, "rd");
1470 break;
1471 case nir_rounding_mode_rtz:
1472 fprintf(fp, "rtz");
1473 break;
1474 default:
1475 fprintf(fp, "unknown");
1476 break;
1477 }
1478 break;
1479 }
1480
1481 case NIR_INTRINSIC_RAY_QUERY_VALUE: {
1482 fprintf(fp, "ray_query_value=");
1483 switch (nir_intrinsic_ray_query_value(instr)) {
1484 #define VAL(_name) \
1485 case nir_ray_query_value_##_name: \
1486 fprintf(fp, #_name); \
1487 break
1488 VAL(intersection_type);
1489 VAL(intersection_t);
1490 VAL(intersection_instance_custom_index);
1491 VAL(intersection_instance_id);
1492 VAL(intersection_instance_sbt_index);
1493 VAL(intersection_geometry_index);
1494 VAL(intersection_primitive_index);
1495 VAL(intersection_barycentrics);
1496 VAL(intersection_front_face);
1497 VAL(intersection_object_ray_direction);
1498 VAL(intersection_object_ray_origin);
1499 VAL(intersection_object_to_world);
1500 VAL(intersection_world_to_object);
1501 VAL(intersection_candidate_aabb_opaque);
1502 VAL(tmin);
1503 VAL(flags);
1504 VAL(world_ray_direction);
1505 VAL(world_ray_origin);
1506 #undef VAL
1507 default:
1508 fprintf(fp, "unknown");
1509 break;
1510 }
1511 break;
1512 }
1513
1514 case NIR_INTRINSIC_RESOURCE_ACCESS_INTEL: {
1515 fprintf(fp, "resource_intel=");
1516 unsigned int modes = nir_intrinsic_resource_access_intel(instr);
1517 if (modes == 0)
1518 fputc('0', fp);
1519 while (modes) {
1520 nir_resource_data_intel i = 1u << u_bit_scan(&modes);
1521 switch (i) {
1522 case nir_resource_intel_bindless:
1523 fprintf(fp, "bindless");
1524 break;
1525 case nir_resource_intel_pushable:
1526 fprintf(fp, "pushable");
1527 break;
1528 case nir_resource_intel_sampler:
1529 fprintf(fp, "sampler");
1530 break;
1531 case nir_resource_intel_non_uniform:
1532 fprintf(fp, "non-uniform");
1533 break;
1534 case nir_resource_intel_sampler_embedded:
1535 fprintf(fp, "sampler-embedded");
1536 break;
1537 default:
1538 fprintf(fp, "unknown");
1539 break;
1540 }
1541 fprintf(fp, "%s", modes ? "|" : "");
1542 }
1543 break;
1544 }
1545
1546 case NIR_INTRINSIC_ACCESS: {
1547 fprintf(fp, "access=");
1548 print_access(nir_intrinsic_access(instr), state, "|");
1549 break;
1550 }
1551
1552 case NIR_INTRINSIC_MATRIX_LAYOUT: {
1553 fprintf(fp, "matrix_layout=");
1554 switch (nir_intrinsic_matrix_layout(instr)) {
1555 case GLSL_MATRIX_LAYOUT_ROW_MAJOR:
1556 fprintf(fp, "row_major");
1557 break;
1558 case GLSL_MATRIX_LAYOUT_COLUMN_MAJOR:
1559 fprintf(fp, "col_major");
1560 break;
1561 default:
1562 fprintf(fp, "unknown");
1563 break;
1564 }
1565 break;
1566 }
1567
1568 case NIR_INTRINSIC_CMAT_DESC: {
1569 struct glsl_cmat_description desc = nir_intrinsic_cmat_desc(instr);
1570 const struct glsl_type *t = glsl_cmat_type(&desc);
1571 fprintf(fp, "%s", glsl_get_type_name(t));
1572 break;
1573 }
1574
1575 case NIR_INTRINSIC_CMAT_SIGNED_MASK: {
1576 fprintf(fp, "cmat_signed=");
1577 unsigned int mask = nir_intrinsic_cmat_signed_mask(instr);
1578 if (mask == 0)
1579 fputc('0', fp);
1580 while (mask) {
1581 nir_cmat_signed i = 1u << u_bit_scan(&mask);
1582 switch (i) {
1583 case NIR_CMAT_A_SIGNED:
1584 fputc('A', fp);
1585 break;
1586 case NIR_CMAT_B_SIGNED:
1587 fputc('B', fp);
1588 break;
1589 case NIR_CMAT_C_SIGNED:
1590 fputc('C', fp);
1591 break;
1592 case NIR_CMAT_RESULT_SIGNED:
1593 fprintf(fp, "Result");
1594 break;
1595 default:
1596 fprintf(fp, "unknown");
1597 break;
1598 }
1599 fprintf(fp, "%s", mask ? "|" : "");
1600 }
1601 break;
1602 }
1603
1604 case NIR_INTRINSIC_ALU_OP: {
1605 nir_op alu_op = nir_intrinsic_alu_op(instr);
1606 fprintf(fp, "alu_op=%s", nir_op_infos[alu_op].name);
1607 break;
1608 }
1609
1610 default: {
1611 unsigned off = info->index_map[idx] - 1;
1612 fprintf(fp, "%s=%d", nir_intrinsic_index_names[idx], instr->const_index[off]);
1613 break;
1614 }
1615 }
1616 }
1617 if (info->num_indices)
1618 fprintf(fp, ")");
1619
1620 if (!state->shader)
1621 return;
1622
1623 nir_variable_mode var_mode;
1624 switch (instr->intrinsic) {
1625 case nir_intrinsic_load_uniform:
1626 var_mode = nir_var_uniform;
1627 break;
1628 case nir_intrinsic_load_input:
1629 case nir_intrinsic_load_per_primitive_input:
1630 case nir_intrinsic_load_interpolated_input:
1631 case nir_intrinsic_load_per_vertex_input:
1632 var_mode = nir_var_shader_in;
1633 break;
1634 case nir_intrinsic_load_output:
1635 case nir_intrinsic_store_output:
1636 case nir_intrinsic_store_per_vertex_output:
1637 var_mode = nir_var_shader_out;
1638 break;
1639 default:
1640 return;
1641 }
1642
1643 if (instr->name) {
1644 fprintf(fp, " // %s", instr->name);
1645 return;
1646 }
1647
1648 nir_foreach_variable_with_modes(var, state->shader, var_mode) {
1649 if (!var->name)
1650 continue;
1651
1652 bool match;
1653 if (instr->intrinsic == nir_intrinsic_load_uniform) {
1654 match = var->data.driver_location == nir_intrinsic_base(instr);
1655 } else {
1656 match = nir_intrinsic_component(instr) >= var->data.location_frac &&
1657 nir_intrinsic_component(instr) <
1658 (var->data.location_frac + glsl_get_components(var->type));
1659 }
1660
1661 if (match) {
1662 fprintf(fp, " // %s", var->name);
1663 break;
1664 }
1665 }
1666 }
1667
1668 static void
print_tex_instr(nir_tex_instr * instr,print_state * state)1669 print_tex_instr(nir_tex_instr *instr, print_state *state)
1670 {
1671 FILE *fp = state->fp;
1672
1673 print_def(&instr->def, state);
1674
1675 fprintf(fp, " = (");
1676 print_alu_type(instr->dest_type, state);
1677 fprintf(fp, ")");
1678
1679 switch (instr->op) {
1680 case nir_texop_tex:
1681 fprintf(fp, "tex ");
1682 break;
1683 case nir_texop_txb:
1684 fprintf(fp, "txb ");
1685 break;
1686 case nir_texop_txl:
1687 fprintf(fp, "txl ");
1688 break;
1689 case nir_texop_txd:
1690 fprintf(fp, "txd ");
1691 break;
1692 case nir_texop_txf:
1693 fprintf(fp, "txf ");
1694 break;
1695 case nir_texop_txf_ms:
1696 fprintf(fp, "txf_ms ");
1697 break;
1698 case nir_texop_txf_ms_fb:
1699 fprintf(fp, "txf_ms_fb ");
1700 break;
1701 case nir_texop_txf_ms_mcs_intel:
1702 fprintf(fp, "txf_ms_mcs_intel ");
1703 break;
1704 case nir_texop_txs:
1705 fprintf(fp, "txs ");
1706 break;
1707 case nir_texop_lod:
1708 fprintf(fp, "lod ");
1709 break;
1710 case nir_texop_tg4:
1711 fprintf(fp, "tg4 ");
1712 break;
1713 case nir_texop_query_levels:
1714 fprintf(fp, "query_levels ");
1715 break;
1716 case nir_texop_texture_samples:
1717 fprintf(fp, "texture_samples ");
1718 break;
1719 case nir_texop_samples_identical:
1720 fprintf(fp, "samples_identical ");
1721 break;
1722 case nir_texop_tex_prefetch:
1723 fprintf(fp, "tex (pre-dispatchable) ");
1724 break;
1725 case nir_texop_fragment_fetch_amd:
1726 fprintf(fp, "fragment_fetch_amd ");
1727 break;
1728 case nir_texop_fragment_mask_fetch_amd:
1729 fprintf(fp, "fragment_mask_fetch_amd ");
1730 break;
1731 case nir_texop_descriptor_amd:
1732 fprintf(fp, "descriptor_amd ");
1733 break;
1734 case nir_texop_sampler_descriptor_amd:
1735 fprintf(fp, "sampler_descriptor_amd ");
1736 break;
1737 case nir_texop_lod_bias_agx:
1738 fprintf(fp, "lod_bias_agx ");
1739 break;
1740 case nir_texop_has_custom_border_color_agx:
1741 fprintf(fp, "has_custom_border_color_agx ");
1742 break;
1743 case nir_texop_custom_border_color_agx:
1744 fprintf(fp, "custom_border_color_agx ");
1745 break;
1746 case nir_texop_hdr_dim_nv:
1747 fprintf(fp, "hdr_dim_nv ");
1748 break;
1749 case nir_texop_tex_type_nv:
1750 fprintf(fp, "tex_type_nv ");
1751 break;
1752 default:
1753 unreachable("Invalid texture operation");
1754 break;
1755 }
1756
1757 bool has_texture_deref = false, has_sampler_deref = false;
1758 for (unsigned i = 0; i < instr->num_srcs; i++) {
1759 if (i > 0) {
1760 fprintf(fp, ", ");
1761 }
1762
1763 print_src(&instr->src[i].src, state, nir_tex_instr_src_type(instr, i));
1764 fprintf(fp, " ");
1765
1766 switch (instr->src[i].src_type) {
1767 case nir_tex_src_backend1:
1768 fprintf(fp, "(backend1)");
1769 break;
1770 case nir_tex_src_backend2:
1771 fprintf(fp, "(backend2)");
1772 break;
1773 case nir_tex_src_coord:
1774 fprintf(fp, "(coord)");
1775 break;
1776 case nir_tex_src_projector:
1777 fprintf(fp, "(projector)");
1778 break;
1779 case nir_tex_src_comparator:
1780 fprintf(fp, "(comparator)");
1781 break;
1782 case nir_tex_src_offset:
1783 fprintf(fp, "(offset)");
1784 break;
1785 case nir_tex_src_bias:
1786 fprintf(fp, "(bias)");
1787 break;
1788 case nir_tex_src_lod:
1789 fprintf(fp, "(lod)");
1790 break;
1791 case nir_tex_src_min_lod:
1792 fprintf(fp, "(min_lod)");
1793 break;
1794 case nir_tex_src_ms_index:
1795 fprintf(fp, "(ms_index)");
1796 break;
1797 case nir_tex_src_ms_mcs_intel:
1798 fprintf(fp, "(ms_mcs_intel)");
1799 break;
1800 case nir_tex_src_ddx:
1801 fprintf(fp, "(ddx)");
1802 break;
1803 case nir_tex_src_ddy:
1804 fprintf(fp, "(ddy)");
1805 break;
1806 case nir_tex_src_sampler_deref_intrinsic:
1807 has_texture_deref = true;
1808 fprintf(fp, "(sampler_deref_intrinsic)");
1809 break;
1810 case nir_tex_src_texture_deref_intrinsic:
1811 has_texture_deref = true;
1812 fprintf(fp, "(texture_deref_intrinsic)");
1813 break;
1814 case nir_tex_src_texture_deref:
1815 has_texture_deref = true;
1816 fprintf(fp, "(texture_deref)");
1817 break;
1818 case nir_tex_src_sampler_deref:
1819 has_sampler_deref = true;
1820 fprintf(fp, "(sampler_deref)");
1821 break;
1822 case nir_tex_src_texture_offset:
1823 fprintf(fp, "(texture_offset)");
1824 break;
1825 case nir_tex_src_sampler_offset:
1826 fprintf(fp, "(sampler_offset)");
1827 break;
1828 case nir_tex_src_texture_handle:
1829 fprintf(fp, "(texture_handle)");
1830 break;
1831 case nir_tex_src_sampler_handle:
1832 fprintf(fp, "(sampler_handle)");
1833 break;
1834 case nir_tex_src_plane:
1835 fprintf(fp, "(plane)");
1836 break;
1837
1838 default:
1839 unreachable("Invalid texture source type");
1840 break;
1841 }
1842 }
1843
1844 if (instr->is_gather_implicit_lod)
1845 fprintf(fp, ", implicit lod");
1846
1847 if (instr->op == nir_texop_tg4) {
1848 fprintf(fp, ", %u (gather_component)", instr->component);
1849 }
1850
1851 if (nir_tex_instr_has_explicit_tg4_offsets(instr)) {
1852 fprintf(fp, ", { (%i, %i)", instr->tg4_offsets[0][0], instr->tg4_offsets[0][1]);
1853 for (unsigned i = 1; i < 4; ++i)
1854 fprintf(fp, ", (%i, %i)", instr->tg4_offsets[i][0],
1855 instr->tg4_offsets[i][1]);
1856 fprintf(fp, " } (offsets)");
1857 }
1858
1859 if (instr->op != nir_texop_txf_ms_fb && !has_texture_deref) {
1860 fprintf(fp, ", %u (texture)", instr->texture_index);
1861 }
1862
1863 if (nir_tex_instr_need_sampler(instr) && !has_sampler_deref) {
1864 fprintf(fp, ", %u (sampler)", instr->sampler_index);
1865 }
1866
1867 if (instr->texture_non_uniform) {
1868 fprintf(fp, ", texture non-uniform");
1869 }
1870
1871 if (instr->sampler_non_uniform) {
1872 fprintf(fp, ", sampler non-uniform");
1873 }
1874
1875 if (instr->is_sparse) {
1876 fprintf(fp, ", sparse");
1877 }
1878 }
1879
1880 static void
print_call_instr(nir_call_instr * instr,print_state * state)1881 print_call_instr(nir_call_instr *instr, print_state *state)
1882 {
1883 FILE *fp = state->fp;
1884
1885 print_no_dest_padding(state);
1886
1887 fprintf(fp, "call %s ", instr->callee->name);
1888
1889 for (unsigned i = 0; i < instr->num_params; i++) {
1890 if (i != 0)
1891 fprintf(fp, ", ");
1892
1893 print_src(&instr->params[i], state, nir_type_invalid);
1894 }
1895 }
1896
1897 static void
print_jump_instr(nir_jump_instr * instr,print_state * state)1898 print_jump_instr(nir_jump_instr *instr, print_state *state)
1899 {
1900 FILE *fp = state->fp;
1901
1902 print_no_dest_padding(state);
1903
1904 switch (instr->type) {
1905 case nir_jump_break:
1906 fprintf(fp, "break");
1907 break;
1908
1909 case nir_jump_continue:
1910 fprintf(fp, "continue");
1911 break;
1912
1913 case nir_jump_return:
1914 fprintf(fp, "return");
1915 break;
1916
1917 case nir_jump_halt:
1918 fprintf(fp, "halt");
1919 break;
1920
1921 case nir_jump_goto:
1922 fprintf(fp, "goto b%u",
1923 instr->target ? instr->target->index : -1);
1924 break;
1925
1926 case nir_jump_goto_if:
1927 fprintf(fp, "goto b%u if ",
1928 instr->target ? instr->target->index : -1);
1929 print_src(&instr->condition, state, nir_type_invalid);
1930 fprintf(fp, " else b%u",
1931 instr->else_target ? instr->else_target->index : -1);
1932 break;
1933 }
1934 }
1935
1936 static void
print_ssa_undef_instr(nir_undef_instr * instr,print_state * state)1937 print_ssa_undef_instr(nir_undef_instr *instr, print_state *state)
1938 {
1939 FILE *fp = state->fp;
1940 print_def(&instr->def, state);
1941 fprintf(fp, " = undefined");
1942 }
1943
1944 static void
print_phi_instr(nir_phi_instr * instr,print_state * state)1945 print_phi_instr(nir_phi_instr *instr, print_state *state)
1946 {
1947 FILE *fp = state->fp;
1948 print_def(&instr->def, state);
1949 fprintf(fp, " = phi ");
1950 nir_foreach_phi_src(src, instr) {
1951 if (&src->node != exec_list_get_head(&instr->srcs))
1952 fprintf(fp, ", ");
1953
1954 fprintf(fp, "b%u: ", src->pred->index);
1955 print_src(&src->src, state, nir_type_invalid);
1956 }
1957 }
1958
1959 static void
print_parallel_copy_instr(nir_parallel_copy_instr * instr,print_state * state)1960 print_parallel_copy_instr(nir_parallel_copy_instr *instr, print_state *state)
1961 {
1962 FILE *fp = state->fp;
1963 nir_foreach_parallel_copy_entry(entry, instr) {
1964 if (&entry->node != exec_list_get_head(&instr->entries))
1965 fprintf(fp, "; ");
1966
1967 if (entry->dest_is_reg) {
1968 fprintf(fp, "*");
1969 print_src(&entry->dest.reg, state, nir_type_invalid);
1970 } else {
1971 print_def(&entry->dest.def, state);
1972 }
1973 fprintf(fp, " = ");
1974
1975 if (entry->src_is_reg)
1976 fprintf(fp, "*");
1977 print_src(&entry->src, state, nir_type_invalid);
1978 }
1979 }
1980
1981 static void
print_debug_info_instr(nir_debug_info_instr * instr,print_state * state)1982 print_debug_info_instr(nir_debug_info_instr *instr, print_state *state)
1983 {
1984 FILE *fp = state->fp;
1985
1986 switch (instr->type) {
1987 case nir_debug_info_src_loc:
1988 fprintf(fp, "// 0x%x", instr->src_loc.spirv_offset);
1989 if (instr->src_loc.line)
1990 fprintf(fp, " %s:%u:%u", nir_src_as_string(instr->src_loc.filename), instr->src_loc.line, instr->src_loc.column);
1991 return;
1992 case nir_debug_info_string:
1993 return; /* Strings are printed for their uses. */
1994 }
1995
1996 unreachable("Unimplemented nir_debug_info_type");
1997 }
1998
1999 static void
print_instr(const nir_instr * instr,print_state * state,unsigned tabs)2000 print_instr(const nir_instr *instr, print_state *state, unsigned tabs)
2001 {
2002 FILE *fp = state->fp;
2003
2004 if (state->debug_info) {
2005 nir_debug_info_instr *di = state->debug_info[instr->index];
2006 di->src_loc.column = (uint32_t)ftell(fp);
2007 }
2008
2009 print_indentation(tabs, fp);
2010
2011 switch (instr->type) {
2012 case nir_instr_type_alu:
2013 print_alu_instr(nir_instr_as_alu(instr), state);
2014 break;
2015
2016 case nir_instr_type_deref:
2017 print_deref_instr(nir_instr_as_deref(instr), state);
2018 break;
2019
2020 case nir_instr_type_call:
2021 print_call_instr(nir_instr_as_call(instr), state);
2022 break;
2023
2024 case nir_instr_type_intrinsic:
2025 print_intrinsic_instr(nir_instr_as_intrinsic(instr), state);
2026 break;
2027
2028 case nir_instr_type_tex:
2029 print_tex_instr(nir_instr_as_tex(instr), state);
2030 break;
2031
2032 case nir_instr_type_load_const:
2033 print_load_const_instr(nir_instr_as_load_const(instr), state);
2034 break;
2035
2036 case nir_instr_type_jump:
2037 print_jump_instr(nir_instr_as_jump(instr), state);
2038 break;
2039
2040 case nir_instr_type_undef:
2041 print_ssa_undef_instr(nir_instr_as_undef(instr), state);
2042 break;
2043
2044 case nir_instr_type_phi:
2045 print_phi_instr(nir_instr_as_phi(instr), state);
2046 break;
2047
2048 case nir_instr_type_parallel_copy:
2049 print_parallel_copy_instr(nir_instr_as_parallel_copy(instr), state);
2050 break;
2051
2052 case nir_instr_type_debug_info:
2053 print_debug_info_instr(nir_instr_as_debug_info(instr), state);
2054 break;
2055
2056 default:
2057 unreachable("Invalid instruction type");
2058 break;
2059 }
2060
2061 if (NIR_DEBUG(PRINT_PASS_FLAGS) && instr->pass_flags)
2062 fprintf(fp, " (pass_flags: 0x%x)", instr->pass_flags);
2063 }
2064
2065 static bool
block_has_instruction_with_dest(nir_block * block)2066 block_has_instruction_with_dest(nir_block *block)
2067 {
2068 nir_foreach_instr(instr, block) {
2069 switch (instr->type) {
2070 case nir_instr_type_load_const:
2071 case nir_instr_type_deref:
2072 case nir_instr_type_alu:
2073 case nir_instr_type_tex:
2074 case nir_instr_type_undef:
2075 case nir_instr_type_phi:
2076 case nir_instr_type_parallel_copy:
2077 return true;
2078
2079 case nir_instr_type_intrinsic: {
2080 nir_intrinsic_instr *intrin = nir_instr_as_intrinsic(instr);
2081 const nir_intrinsic_info *info = &nir_intrinsic_infos[intrin->intrinsic];
2082 if (info->has_dest)
2083 return true;
2084
2085 /* Doesn't define a new value. */
2086 break;
2087 }
2088
2089 case nir_instr_type_jump:
2090 case nir_instr_type_call:
2091 case nir_instr_type_debug_info:
2092 /* Doesn't define a new value. */
2093 break;
2094 }
2095 }
2096
2097 return false;
2098 }
2099
2100 static void print_cf_node(nir_cf_node *node, print_state *state,
2101 unsigned tabs);
2102
2103 static void
print_block_preds(nir_block * block,print_state * state)2104 print_block_preds(nir_block *block, print_state *state)
2105 {
2106 FILE *fp = state->fp;
2107 nir_block **preds = nir_block_get_predecessors_sorted(block, NULL);
2108 for (unsigned i = 0; i < block->predecessors->entries; i++) {
2109 if (i != 0)
2110 fprintf(fp, " ");
2111 fprintf(fp, "b%u", preds[i]->index);
2112 }
2113 ralloc_free(preds);
2114 }
2115
2116 static void
print_block_succs(nir_block * block,print_state * state)2117 print_block_succs(nir_block *block, print_state *state)
2118 {
2119 FILE *fp = state->fp;
2120 for (unsigned i = 0; i < 2; i++) {
2121 if (block->successors[i]) {
2122 fprintf(fp, "b%u ", block->successors[i]->index);
2123 }
2124 }
2125 }
2126
2127 static void
print_block(nir_block * block,print_state * state,unsigned tabs)2128 print_block(nir_block *block, print_state *state, unsigned tabs)
2129 {
2130 FILE *fp = state->fp;
2131
2132 if (block_has_instruction_with_dest(block))
2133 state->padding_for_no_dest = calculate_padding_for_no_dest(state);
2134 else
2135 state->padding_for_no_dest = 0;
2136
2137 print_indentation(tabs, fp);
2138 fprintf(fp, "%sblock b%u:",
2139 divergence_status(state, block->divergent),
2140 block->index);
2141
2142 const bool empty_block = exec_list_is_empty(&block->instr_list);
2143 if (empty_block) {
2144 fprintf(fp, " // preds: ");
2145 print_block_preds(block, state);
2146 fprintf(fp, ", succs: ");
2147 print_block_succs(block, state);
2148 fprintf(fp, "\n");
2149 return;
2150 }
2151
2152 const unsigned block_length = 7 + count_digits(block->index) + 1;
2153 const unsigned pred_padding = block_length < state->padding_for_no_dest ? state->padding_for_no_dest - block_length : 0;
2154
2155 fprintf(fp, "%*s// preds: ", pred_padding, "");
2156 print_block_preds(block, state);
2157 fprintf(fp, "\n");
2158
2159 nir_foreach_instr(instr, block) {
2160 print_instr(instr, state, tabs);
2161 fprintf(fp, "\n");
2162 print_annotation(state, instr);
2163 }
2164
2165 print_indentation(tabs, fp);
2166 fprintf(fp, "%*s// succs: ", state->padding_for_no_dest, "");
2167 print_block_succs(block, state);
2168 fprintf(fp, "\n");
2169 }
2170
2171 static void
print_if(nir_if * if_stmt,print_state * state,unsigned tabs)2172 print_if(nir_if *if_stmt, print_state *state, unsigned tabs)
2173 {
2174 FILE *fp = state->fp;
2175
2176 print_indentation(tabs, fp);
2177 fprintf(fp, "if ");
2178 print_src(&if_stmt->condition, state, nir_type_invalid);
2179 switch (if_stmt->control) {
2180 case nir_selection_control_flatten:
2181 fprintf(fp, " // flatten");
2182 break;
2183 case nir_selection_control_dont_flatten:
2184 fprintf(fp, " // don't flatten");
2185 break;
2186 case nir_selection_control_divergent_always_taken:
2187 fprintf(fp, " // divergent always taken");
2188 break;
2189 case nir_selection_control_none:
2190 default:
2191 break;
2192 }
2193 fprintf(fp, " {\n");
2194 foreach_list_typed(nir_cf_node, node, node, &if_stmt->then_list) {
2195 print_cf_node(node, state, tabs + 1);
2196 }
2197 print_indentation(tabs, fp);
2198 fprintf(fp, "} else {\n");
2199 foreach_list_typed(nir_cf_node, node, node, &if_stmt->else_list) {
2200 print_cf_node(node, state, tabs + 1);
2201 }
2202 print_indentation(tabs, fp);
2203 fprintf(fp, "}\n");
2204 }
2205
2206 static void
print_loop(nir_loop * loop,print_state * state,unsigned tabs)2207 print_loop(nir_loop *loop, print_state *state, unsigned tabs)
2208 {
2209 FILE *fp = state->fp;
2210
2211 print_indentation(tabs, fp);
2212 fprintf(fp, "%sloop {\n", divergence_status(state, loop->divergent));
2213 foreach_list_typed(nir_cf_node, node, node, &loop->body) {
2214 print_cf_node(node, state, tabs + 1);
2215 }
2216 print_indentation(tabs, fp);
2217
2218 if (nir_loop_has_continue_construct(loop)) {
2219 fprintf(fp, "} continue {\n");
2220 foreach_list_typed(nir_cf_node, node, node, &loop->continue_list) {
2221 print_cf_node(node, state, tabs + 1);
2222 }
2223 print_indentation(tabs, fp);
2224 }
2225
2226 fprintf(fp, "}\n");
2227 }
2228
2229 static void
print_cf_node(nir_cf_node * node,print_state * state,unsigned int tabs)2230 print_cf_node(nir_cf_node *node, print_state *state, unsigned int tabs)
2231 {
2232 switch (node->type) {
2233 case nir_cf_node_block:
2234 print_block(nir_cf_node_as_block(node), state, tabs);
2235 break;
2236
2237 case nir_cf_node_if:
2238 print_if(nir_cf_node_as_if(node), state, tabs);
2239 break;
2240
2241 case nir_cf_node_loop:
2242 print_loop(nir_cf_node_as_loop(node), state, tabs);
2243 break;
2244
2245 default:
2246 unreachable("Invalid CFG node type");
2247 }
2248 }
2249
2250 static void
print_function_impl(nir_function_impl * impl,print_state * state)2251 print_function_impl(nir_function_impl *impl, print_state *state)
2252 {
2253 FILE *fp = state->fp;
2254
2255 state->max_dest_index = impl->ssa_alloc;
2256
2257 fprintf(fp, "\nimpl %s ", impl->function->name);
2258
2259 fprintf(fp, "{\n");
2260
2261 if (impl->preamble) {
2262 print_indentation(1, fp);
2263 fprintf(fp, "preamble %s\n", impl->preamble->name);
2264 }
2265
2266 if (!NIR_DEBUG(PRINT_NO_INLINE_CONSTS)) {
2267 /* Don't reindex the SSA as suggested by nir_gather_types() because
2268 * nir_print don't modify the shader. If needed, a limit for ssa_alloc
2269 * can be added.
2270 */
2271 state->float_types = calloc(BITSET_WORDS(impl->ssa_alloc), sizeof(BITSET_WORD));
2272 state->int_types = calloc(BITSET_WORDS(impl->ssa_alloc), sizeof(BITSET_WORD));
2273 nir_gather_types(impl, state->float_types, state->int_types);
2274 }
2275
2276 nir_foreach_function_temp_variable(var, impl) {
2277 print_indentation(1, fp);
2278 print_var_decl(var, state);
2279 }
2280
2281 nir_index_blocks(impl);
2282
2283 foreach_list_typed(nir_cf_node, node, node, &impl->body) {
2284 print_cf_node(node, state, 1);
2285 }
2286
2287 print_indentation(1, fp);
2288 fprintf(fp, "block b%u:\n}\n\n", impl->end_block->index);
2289
2290 free(state->float_types);
2291 free(state->int_types);
2292 state->max_dest_index = 0;
2293 }
2294
2295 static void
print_function(nir_function * function,print_state * state)2296 print_function(nir_function *function, print_state *state)
2297 {
2298 FILE *fp = state->fp;
2299
2300 /* clang-format off */
2301 fprintf(fp, "decl_function %s (%d params)%s%s", function->name,
2302 function->num_params,
2303 function->dont_inline ? " (noinline)" :
2304 function->should_inline ? " (inline)" : "",
2305 function->is_exported ? " (exported)" : "");
2306 /* clang-format on */
2307
2308 fprintf(fp, "\n");
2309
2310 if (function->impl != NULL) {
2311 print_function_impl(function->impl, state);
2312 return;
2313 }
2314 }
2315
2316 static void
init_print_state(print_state * state,nir_shader * shader,FILE * fp)2317 init_print_state(print_state *state, nir_shader *shader, FILE *fp)
2318 {
2319 state->fp = fp;
2320 state->shader = shader;
2321 state->ht = _mesa_pointer_hash_table_create(NULL);
2322 state->syms = _mesa_set_create(NULL, _mesa_hash_string,
2323 _mesa_key_string_equal);
2324 state->index = 0;
2325 state->int_types = NULL;
2326 state->float_types = NULL;
2327 state->max_dest_index = 0;
2328 state->padding_for_no_dest = 0;
2329 }
2330
2331 static void
destroy_print_state(print_state * state)2332 destroy_print_state(print_state *state)
2333 {
2334 _mesa_hash_table_destroy(state->ht, NULL);
2335 _mesa_set_destroy(state->syms, NULL);
2336 }
2337
2338 static const char *
primitive_name(unsigned primitive)2339 primitive_name(unsigned primitive)
2340 {
2341 #define PRIM(X) \
2342 case MESA_PRIM_##X: \
2343 return #X
2344 switch (primitive) {
2345 PRIM(POINTS);
2346 PRIM(LINES);
2347 PRIM(LINE_LOOP);
2348 PRIM(LINE_STRIP);
2349 PRIM(TRIANGLES);
2350 PRIM(TRIANGLE_STRIP);
2351 PRIM(TRIANGLE_FAN);
2352 PRIM(QUADS);
2353 PRIM(QUAD_STRIP);
2354 PRIM(POLYGON);
2355 PRIM(LINES_ADJACENCY);
2356 PRIM(TRIANGLES_ADJACENCY);
2357 default:
2358 return "UNKNOWN";
2359 }
2360 }
2361
2362 static void
print_bitset(FILE * fp,const char * label,const unsigned * words,int size)2363 print_bitset(FILE *fp, const char *label, const unsigned *words, int size)
2364 {
2365 fprintf(fp, "%s: ", label);
2366 /* Iterate back-to-front to get proper digit order (most significant first). */
2367 for (int i = size - 1; i >= 0; --i) {
2368 fprintf(fp, (i == size - 1) ? "0x%08x" : "'%08x", words[i]);
2369 }
2370 fprintf(fp, "\n");
2371 }
2372
2373 /* Print bitset, only if some bits are set */
2374 static void
print_nz_bitset(FILE * fp,const char * label,const unsigned * words,int size)2375 print_nz_bitset(FILE *fp, const char *label, const unsigned *words, int size)
2376 {
2377 bool is_all_zero = true;
2378 for (int i = 0; i < size; ++i) {
2379 if (words[i]) {
2380 is_all_zero = false;
2381 break;
2382 }
2383 }
2384
2385 if (!is_all_zero)
2386 print_bitset(fp, label, words, size);
2387 }
2388
2389 /* Print uint64_t value, only if non-zero.
2390 * The value is printed by enumerating the ranges of bits that are set.
2391 * E.g. inputs_read: 0,15-17
2392 */
2393 static void
print_nz_x64(FILE * fp,const char * label,uint64_t value)2394 print_nz_x64(FILE *fp, const char *label, uint64_t value)
2395 {
2396 if (value) {
2397 char acc[256] = { 0 };
2398 char buf[32];
2399 int start = 0;
2400 int count = 0;
2401 while (value) {
2402 u_bit_scan_consecutive_range64(&value, &start, &count);
2403 assert(count > 0);
2404 bool is_first = !acc[0];
2405 if (count > 1) {
2406 snprintf(buf, sizeof(buf), is_first ? "%d-%d" : ",%d-%d", start, start + count - 1);
2407 } else {
2408 snprintf(buf, sizeof(buf), is_first ? "%d" : ",%d", start);
2409 }
2410 assert(strlen(acc) + strlen(buf) + 1 < sizeof(acc));
2411 strcat(acc, buf);
2412 }
2413 fprintf(fp, "%s: %s\n", label, acc);
2414 }
2415 }
2416
2417 /* Print uint32_t value in hex, only if non-zero */
2418 static void
print_nz_x32(FILE * fp,const char * label,uint32_t value)2419 print_nz_x32(FILE *fp, const char *label, uint32_t value)
2420 {
2421 if (value)
2422 fprintf(fp, "%s: 0x%08" PRIx32 "\n", label, value);
2423 }
2424
2425 /* Print uint16_t value in hex, only if non-zero */
2426 static void
print_nz_x16(FILE * fp,const char * label,uint16_t value)2427 print_nz_x16(FILE *fp, const char *label, uint16_t value)
2428 {
2429 if (value)
2430 fprintf(fp, "%s: 0x%04x\n", label, value);
2431 }
2432
2433 /* Print uint8_t value in hex, only if non-zero */
2434 static void
print_nz_x8(FILE * fp,const char * label,uint8_t value)2435 print_nz_x8(FILE *fp, const char *label, uint8_t value)
2436 {
2437 if (value)
2438 fprintf(fp, "%s: 0x%02x\n", label, value);
2439 }
2440
2441 /* Print unsigned value in decimal, only if non-zero */
2442 static void
print_nz_unsigned(FILE * fp,const char * label,unsigned value)2443 print_nz_unsigned(FILE *fp, const char *label, unsigned value)
2444 {
2445 if (value)
2446 fprintf(fp, "%s: %u\n", label, value);
2447 }
2448
2449 /* Print bool only if set */
2450 static void
print_nz_bool(FILE * fp,const char * label,bool value)2451 print_nz_bool(FILE *fp, const char *label, bool value)
2452 {
2453 if (value)
2454 fprintf(fp, "%s: true\n", label);
2455 }
2456
2457 static void
print_shader_info(const struct shader_info * info,FILE * fp)2458 print_shader_info(const struct shader_info *info, FILE *fp)
2459 {
2460 fprintf(fp, "shader: %s\n", gl_shader_stage_name(info->stage));
2461
2462 fprintf(fp, "source_blake3: {");
2463 _mesa_blake3_print(fp, info->source_blake3);
2464 fprintf(fp, "}\n");
2465
2466 if (info->name)
2467 fprintf(fp, "name: %s\n", info->name);
2468
2469 if (info->label)
2470 fprintf(fp, "label: %s\n", info->label);
2471
2472 fprintf(fp, "internal: %s\n", info->internal ? "true" : "false");
2473
2474 if (gl_shader_stage_uses_workgroup(info->stage)) {
2475 fprintf(fp, "workgroup_size: %u, %u, %u%s\n",
2476 info->workgroup_size[0],
2477 info->workgroup_size[1],
2478 info->workgroup_size[2],
2479 info->workgroup_size_variable ? " (variable)" : "");
2480 }
2481
2482 fprintf(fp, "stage: %d\n"
2483 "next_stage: %d\n",
2484 info->stage, info->next_stage);
2485
2486 print_nz_unsigned(fp, "num_textures", info->num_textures);
2487 print_nz_unsigned(fp, "num_ubos", info->num_ubos);
2488 print_nz_unsigned(fp, "num_abos", info->num_abos);
2489 print_nz_unsigned(fp, "num_ssbos", info->num_ssbos);
2490 print_nz_unsigned(fp, "num_images", info->num_images);
2491
2492 print_nz_x64(fp, "inputs_read", info->inputs_read);
2493 print_nz_x64(fp, "dual_slot_inputs", info->dual_slot_inputs);
2494 print_nz_x64(fp, "outputs_written", info->outputs_written);
2495 print_nz_x64(fp, "outputs_read", info->outputs_read);
2496
2497 print_nz_bitset(fp, "system_values_read", info->system_values_read, ARRAY_SIZE(info->system_values_read));
2498
2499 print_nz_x64(fp, "per_primitive_inputs", info->per_primitive_inputs);
2500 print_nz_x64(fp, "per_primitive_outputs", info->per_primitive_outputs);
2501 print_nz_x64(fp, "per_view_outputs", info->per_view_outputs);
2502
2503 print_nz_x16(fp, "inputs_read_16bit", info->inputs_read_16bit);
2504 print_nz_x16(fp, "outputs_written_16bit", info->outputs_written_16bit);
2505 print_nz_x16(fp, "outputs_read_16bit", info->outputs_read_16bit);
2506 print_nz_x16(fp, "inputs_read_indirectly_16bit", info->inputs_read_indirectly_16bit);
2507 print_nz_x16(fp, "outputs_accessed_indirectly_16bit", info->outputs_accessed_indirectly_16bit);
2508
2509 print_nz_x32(fp, "patch_inputs_read", info->patch_inputs_read);
2510 print_nz_x32(fp, "patch_outputs_written", info->patch_outputs_written);
2511 print_nz_x32(fp, "patch_outputs_read", info->patch_outputs_read);
2512
2513 print_nz_x64(fp, "inputs_read_indirectly", info->inputs_read_indirectly);
2514 print_nz_x64(fp, "outputs_accessed_indirectly", info->outputs_accessed_indirectly);
2515 print_nz_x64(fp, "patch_inputs_read_indirectly", info->patch_inputs_read_indirectly);
2516 print_nz_x64(fp, "patch_outputs_accessed_indirectly", info->patch_outputs_accessed_indirectly);
2517
2518 print_nz_bitset(fp, "textures_used", info->textures_used, ARRAY_SIZE(info->textures_used));
2519 print_nz_bitset(fp, "textures_used_by_txf", info->textures_used_by_txf, ARRAY_SIZE(info->textures_used_by_txf));
2520 print_nz_bitset(fp, "samplers_used", info->samplers_used, ARRAY_SIZE(info->samplers_used));
2521 print_nz_bitset(fp, "images_used", info->images_used, ARRAY_SIZE(info->images_used));
2522 print_nz_bitset(fp, "image_buffers", info->image_buffers, ARRAY_SIZE(info->image_buffers));
2523 print_nz_bitset(fp, "msaa_images", info->msaa_images, ARRAY_SIZE(info->msaa_images));
2524
2525 print_nz_x32(fp, "float_controls_execution_mode", info->float_controls_execution_mode);
2526
2527 print_nz_unsigned(fp, "shared_size", info->shared_size);
2528
2529 if (info->stage == MESA_SHADER_MESH || info->stage == MESA_SHADER_TASK) {
2530 fprintf(fp, "task_payload_size: %u\n", info->task_payload_size);
2531 }
2532
2533 print_nz_unsigned(fp, "ray queries", info->ray_queries);
2534
2535 fprintf(fp, "subgroup_size: %u\n", info->subgroup_size);
2536
2537 print_nz_bool(fp, "uses_wide_subgroup_intrinsics", info->uses_wide_subgroup_intrinsics);
2538
2539 bool has_xfb_stride = info->xfb_stride[0] || info->xfb_stride[1] || info->xfb_stride[2] || info->xfb_stride[3];
2540 if (has_xfb_stride)
2541 fprintf(fp, "xfb_stride: {%u, %u, %u, %u}\n",
2542 info->xfb_stride[0],
2543 info->xfb_stride[1],
2544 info->xfb_stride[2],
2545 info->xfb_stride[3]);
2546
2547 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];
2548 if (has_inlinable_uniform_dw_offsets)
2549 fprintf(fp, "inlinable_uniform_dw_offsets: {%u, %u, %u, %u}\n",
2550 info->inlinable_uniform_dw_offsets[0],
2551 info->inlinable_uniform_dw_offsets[1],
2552 info->inlinable_uniform_dw_offsets[2],
2553 info->inlinable_uniform_dw_offsets[3]);
2554
2555 print_nz_unsigned(fp, "num_inlinable_uniforms", info->num_inlinable_uniforms);
2556 print_nz_unsigned(fp, "clip_distance_array_size", info->clip_distance_array_size);
2557 print_nz_unsigned(fp, "cull_distance_array_size", info->cull_distance_array_size);
2558
2559 print_nz_bool(fp, "uses_texture_gather", info->uses_texture_gather);
2560 print_nz_bool(fp, "uses_resource_info_query", info->uses_resource_info_query);
2561 print_nz_bool(fp, "uses_fddx_fddy", info->uses_fddx_fddy);
2562 print_nz_bool(fp, "divergence_analysis_run", info->divergence_analysis_run);
2563
2564 print_nz_x8(fp, "bit_sizes_float", info->bit_sizes_float);
2565 print_nz_x8(fp, "bit_sizes_int", info->bit_sizes_int);
2566
2567 print_nz_bool(fp, "first_ubo_is_default_ubo", info->first_ubo_is_default_ubo);
2568 print_nz_bool(fp, "separate_shader", info->separate_shader);
2569 print_nz_bool(fp, "has_transform_feedback_varyings", info->has_transform_feedback_varyings);
2570 print_nz_bool(fp, "flrp_lowered", info->flrp_lowered);
2571 print_nz_bool(fp, "io_lowered", info->io_lowered);
2572 print_nz_bool(fp, "writes_memory", info->writes_memory);
2573 print_nz_unsigned(fp, "derivative_group", info->derivative_group);
2574
2575 switch (info->stage) {
2576 case MESA_SHADER_VERTEX:
2577 print_nz_x64(fp, "double_inputs", info->vs.double_inputs);
2578 print_nz_unsigned(fp, "blit_sgprs_amd", info->vs.blit_sgprs_amd);
2579 print_nz_bool(fp, "window_space_position", info->vs.window_space_position);
2580 print_nz_bool(fp, "needs_edge_flag", info->vs.needs_edge_flag);
2581 break;
2582
2583 case MESA_SHADER_TESS_CTRL:
2584 case MESA_SHADER_TESS_EVAL:
2585 fprintf(fp, "primitive_mode: %u\n", info->tess._primitive_mode);
2586 fprintf(fp, "tcs_vertices_out: %u\n", info->tess.tcs_vertices_out);
2587 fprintf(fp, "spacing: %u\n", info->tess.spacing);
2588
2589 print_nz_bool(fp, "ccw", info->tess.ccw);
2590 print_nz_bool(fp, "point_mode", info->tess.point_mode);
2591 print_nz_x64(fp, "tcs_cross_invocation_inputs_read", info->tess.tcs_cross_invocation_inputs_read);
2592 print_nz_x64(fp, "tcs_cross_invocation_outputs_read", info->tess.tcs_cross_invocation_outputs_read);
2593 break;
2594
2595 case MESA_SHADER_GEOMETRY:
2596 fprintf(fp, "output_primitive: %s\n", primitive_name(info->gs.output_primitive));
2597 fprintf(fp, "input_primitive: %s\n", primitive_name(info->gs.input_primitive));
2598 fprintf(fp, "vertices_out: %u\n", info->gs.vertices_out);
2599 fprintf(fp, "invocations: %u\n", info->gs.invocations);
2600 fprintf(fp, "vertices_in: %u\n", info->gs.vertices_in);
2601 print_nz_bool(fp, "uses_end_primitive", info->gs.uses_end_primitive);
2602 fprintf(fp, "active_stream_mask: 0x%02x\n", info->gs.active_stream_mask);
2603 break;
2604
2605 case MESA_SHADER_FRAGMENT:
2606 print_nz_bool(fp, "uses_discard", info->fs.uses_discard);
2607 print_nz_bool(fp, "uses_fbfetch_output", info->fs.uses_fbfetch_output);
2608 print_nz_bool(fp, "color_is_dual_source", info->fs.color_is_dual_source);
2609
2610 print_nz_bool(fp, "require_full_quads", info->fs.require_full_quads);
2611 print_nz_bool(fp, "needs_quad_helper_invocations", info->fs.needs_quad_helper_invocations);
2612 print_nz_bool(fp, "uses_sample_qualifier", info->fs.uses_sample_qualifier);
2613 print_nz_bool(fp, "uses_sample_shading", info->fs.uses_sample_shading);
2614 print_nz_bool(fp, "early_fragment_tests", info->fs.early_fragment_tests);
2615 print_nz_bool(fp, "inner_coverage", info->fs.inner_coverage);
2616 print_nz_bool(fp, "post_depth_coverage", info->fs.post_depth_coverage);
2617
2618 print_nz_bool(fp, "pixel_center_integer", info->fs.pixel_center_integer);
2619 print_nz_bool(fp, "origin_upper_left", info->fs.origin_upper_left);
2620 print_nz_bool(fp, "pixel_interlock_ordered", info->fs.pixel_interlock_ordered);
2621 print_nz_bool(fp, "pixel_interlock_unordered", info->fs.pixel_interlock_unordered);
2622 print_nz_bool(fp, "sample_interlock_ordered", info->fs.sample_interlock_ordered);
2623 print_nz_bool(fp, "sample_interlock_unordered", info->fs.sample_interlock_unordered);
2624 print_nz_bool(fp, "untyped_color_outputs", info->fs.untyped_color_outputs);
2625
2626 print_nz_unsigned(fp, "depth_layout", info->fs.depth_layout);
2627
2628 if (info->fs.color0_interp != INTERP_MODE_NONE) {
2629 fprintf(fp, "color0_interp: %s\n",
2630 glsl_interp_mode_name(info->fs.color0_interp));
2631 }
2632 print_nz_bool(fp, "color0_sample", info->fs.color0_sample);
2633 print_nz_bool(fp, "color0_centroid", info->fs.color0_centroid);
2634
2635 if (info->fs.color1_interp != INTERP_MODE_NONE) {
2636 fprintf(fp, "color1_interp: %s\n",
2637 glsl_interp_mode_name(info->fs.color1_interp));
2638 }
2639 print_nz_bool(fp, "color1_sample", info->fs.color1_sample);
2640 print_nz_bool(fp, "color1_centroid", info->fs.color1_centroid);
2641
2642 print_nz_x32(fp, "advanced_blend_modes", info->fs.advanced_blend_modes);
2643 break;
2644
2645 case MESA_SHADER_COMPUTE:
2646 case MESA_SHADER_KERNEL:
2647 if (info->cs.workgroup_size_hint[0] || info->cs.workgroup_size_hint[1] || info->cs.workgroup_size_hint[2])
2648 fprintf(fp, "workgroup_size_hint: {%u, %u, %u}\n",
2649 info->cs.workgroup_size_hint[0],
2650 info->cs.workgroup_size_hint[1],
2651 info->cs.workgroup_size_hint[2]);
2652 print_nz_unsigned(fp, "user_data_components_amd", info->cs.user_data_components_amd);
2653 fprintf(fp, "ptr_size: %u\n", info->cs.ptr_size);
2654 break;
2655
2656 case MESA_SHADER_MESH:
2657 print_nz_x64(fp, "ms_cross_invocation_output_access", info->mesh.ms_cross_invocation_output_access);
2658 fprintf(fp, "max_vertices_out: %u\n", info->mesh.max_vertices_out);
2659 fprintf(fp, "max_primitives_out: %u\n", info->mesh.max_primitives_out);
2660 fprintf(fp, "primitive_type: %s\n", primitive_name(info->mesh.primitive_type));
2661 print_nz_bool(fp, "nv", info->mesh.nv);
2662 break;
2663
2664 default:
2665 fprintf(fp, "Unhandled stage %d\n", info->stage);
2666 }
2667 }
2668
2669 static void
_nir_print_shader_annotated(nir_shader * shader,FILE * fp,struct hash_table * annotations,nir_debug_info_instr ** debug_info)2670 _nir_print_shader_annotated(nir_shader *shader, FILE *fp,
2671 struct hash_table *annotations,
2672 nir_debug_info_instr **debug_info)
2673 {
2674 print_state state;
2675 init_print_state(&state, shader, fp);
2676 state.def_prefix = debug_info ? "ssa_" : "%";
2677 state.annotations = annotations;
2678 state.debug_info = debug_info;
2679
2680 print_shader_info(&shader->info, fp);
2681
2682 fprintf(fp, "inputs: %u\n", shader->num_inputs);
2683 fprintf(fp, "outputs: %u\n", shader->num_outputs);
2684 fprintf(fp, "uniforms: %u\n", shader->num_uniforms);
2685 if (shader->scratch_size)
2686 fprintf(fp, "scratch: %u\n", shader->scratch_size);
2687 if (shader->constant_data_size)
2688 fprintf(fp, "constants: %u\n", shader->constant_data_size);
2689 for (unsigned i = 0; i < nir_num_variable_modes; i++) {
2690 nir_variable_mode mode = BITFIELD_BIT(i);
2691 if (mode == nir_var_function_temp)
2692 continue;
2693
2694 if (mode == nir_var_shader_in || mode == nir_var_shader_out) {
2695 for (unsigned j = 0; j < 128; j++) {
2696 nir_variable *vars[NIR_MAX_VEC_COMPONENTS] = {0};
2697 nir_foreach_variable_with_modes(var, shader, mode) {
2698 if (var->data.location == j)
2699 vars[var->data.location_frac] = var;
2700 }
2701 for (unsigned j = 0; j < ARRAY_SIZE(vars); j++)
2702 if (vars[j]) {
2703 print_var_decl(vars[j], &state);
2704 }
2705 }
2706 } else {
2707 nir_foreach_variable_with_modes(var, shader, mode)
2708 print_var_decl(var, &state);
2709 }
2710 }
2711
2712 foreach_list_typed(nir_function, func, node, &shader->functions) {
2713 print_function(func, &state);
2714 }
2715
2716 destroy_print_state(&state);
2717 }
2718
2719 void
nir_print_shader_annotated(nir_shader * shader,FILE * fp,struct hash_table * annotations)2720 nir_print_shader_annotated(nir_shader *shader, FILE *fp,
2721 struct hash_table *annotations)
2722 {
2723 _nir_print_shader_annotated(shader, fp, annotations, NULL);
2724 }
2725
2726 void
nir_print_shader(nir_shader * shader,FILE * fp)2727 nir_print_shader(nir_shader *shader, FILE *fp)
2728 {
2729 nir_print_shader_annotated(shader, fp, NULL);
2730 fflush(fp);
2731 }
2732
2733 static char *
_nir_shader_as_str_annotated(nir_shader * nir,struct hash_table * annotations,void * mem_ctx,nir_debug_info_instr ** debug_info)2734 _nir_shader_as_str_annotated(nir_shader *nir, struct hash_table *annotations, void *mem_ctx,
2735 nir_debug_info_instr **debug_info)
2736 {
2737 char *stream_data = NULL;
2738 size_t stream_size = 0;
2739 struct u_memstream mem;
2740 if (u_memstream_open(&mem, &stream_data, &stream_size)) {
2741 FILE *const stream = u_memstream_get(&mem);
2742 _nir_print_shader_annotated(nir, stream, annotations, debug_info);
2743 u_memstream_close(&mem);
2744 }
2745
2746 char *str = ralloc_size(mem_ctx, stream_size + 1);
2747 memcpy(str, stream_data, stream_size);
2748 str[stream_size] = '\0';
2749
2750 free(stream_data);
2751
2752 return str;
2753 }
2754
2755 char *
nir_shader_as_str_annotated(nir_shader * nir,struct hash_table * annotations,void * mem_ctx)2756 nir_shader_as_str_annotated(nir_shader *nir, struct hash_table *annotations, void *mem_ctx)
2757 {
2758 return _nir_shader_as_str_annotated(nir, annotations, mem_ctx, NULL);
2759 }
2760
2761 char *
nir_shader_as_str(nir_shader * nir,void * mem_ctx)2762 nir_shader_as_str(nir_shader *nir, void *mem_ctx)
2763 {
2764 return nir_shader_as_str_annotated(nir, NULL, mem_ctx);
2765 }
2766
2767 void
nir_print_instr(const nir_instr * instr,FILE * fp)2768 nir_print_instr(const nir_instr *instr, FILE *fp)
2769 {
2770 print_state state = {
2771 .fp = fp,
2772 .def_prefix = "%",
2773 };
2774 if (instr->block) {
2775 nir_function_impl *impl = nir_cf_node_get_function(&instr->block->cf_node);
2776 state.shader = impl->function->shader;
2777 }
2778
2779 print_instr(instr, &state, 0);
2780 }
2781
2782 char *
nir_instr_as_str(const nir_instr * instr,void * mem_ctx)2783 nir_instr_as_str(const nir_instr *instr, void *mem_ctx)
2784 {
2785 char *stream_data = NULL;
2786 size_t stream_size = 0;
2787 struct u_memstream mem;
2788 if (u_memstream_open(&mem, &stream_data, &stream_size)) {
2789 FILE *const stream = u_memstream_get(&mem);
2790 nir_print_instr(instr, stream);
2791 u_memstream_close(&mem);
2792 }
2793
2794 char *str = ralloc_size(mem_ctx, stream_size + 1);
2795 memcpy(str, stream_data, stream_size);
2796 str[stream_size] = '\0';
2797
2798 free(stream_data);
2799
2800 return str;
2801 }
2802
2803 void
nir_print_deref(const nir_deref_instr * deref,FILE * fp)2804 nir_print_deref(const nir_deref_instr *deref, FILE *fp)
2805 {
2806 print_state state = {
2807 .fp = fp,
2808 .def_prefix = "%",
2809 };
2810 print_deref_link(deref, true, &state);
2811 }
2812
2813 void
nir_log_shader_annotated_tagged(enum mesa_log_level level,const char * tag,nir_shader * shader,struct hash_table * annotations)2814 nir_log_shader_annotated_tagged(enum mesa_log_level level, const char *tag,
2815 nir_shader *shader, struct hash_table *annotations)
2816 {
2817 char *str = nir_shader_as_str_annotated(shader, annotations, NULL);
2818 _mesa_log_multiline(level, tag, str);
2819 ralloc_free(str);
2820 }
2821
2822 char *
nir_shader_gather_debug_info(nir_shader * shader,const char * filename)2823 nir_shader_gather_debug_info(nir_shader *shader, const char *filename)
2824 {
2825 uint32_t instr_count = 0;
2826 nir_foreach_function_impl(impl, shader) {
2827 nir_foreach_block(block, impl) {
2828 nir_foreach_instr(instr, block) {
2829 instr->index = instr_count;
2830 instr_count++;
2831 }
2832 }
2833 }
2834
2835 if (!instr_count)
2836 return nir_shader_as_str(shader, NULL);
2837
2838 nir_debug_info_instr **debug_info = rzalloc_array(shader, nir_debug_info_instr *, instr_count);
2839
2840 instr_count = 0;
2841 nir_foreach_function_impl(impl, shader) {
2842 nir_builder b = nir_builder_at(nir_before_cf_list(&impl->body));
2843 nir_def *filename_def = nir_build_string(&b, filename);
2844
2845 nir_foreach_block(block, impl) {
2846 nir_foreach_instr_safe(instr, block) {
2847 if (instr->type == nir_instr_type_debug_info)
2848 continue;
2849
2850 nir_debug_info_instr *di = nir_debug_info_instr_create(shader, nir_debug_info_src_loc, 0);
2851 di->src_loc.filename = nir_src_for_ssa(filename_def);
2852 di->src_loc.source = nir_debug_info_nir;
2853 debug_info[instr_count++] = di;
2854 }
2855 }
2856 }
2857
2858 char *str = _nir_shader_as_str_annotated(shader, NULL, NULL, debug_info);
2859
2860 uint32_t line = 1;
2861 uint32_t character_index = 0;
2862
2863 for (uint32_t i = 0; i < instr_count; i++) {
2864 nir_debug_info_instr *di = debug_info[i];
2865
2866 while (character_index < di->src_loc.column) {
2867 if (str[character_index] == '\n')
2868 line++;
2869 character_index++;
2870 }
2871
2872 di->src_loc.line = line;
2873 di->src_loc.column = 0;
2874 }
2875
2876 instr_count = 0;
2877 nir_foreach_function_impl(impl, shader) {
2878 nir_foreach_block(block, impl) {
2879 nir_foreach_instr_safe(instr, block) {
2880 if (instr->type != nir_instr_type_debug_info)
2881 nir_instr_insert_before(instr, &debug_info[instr_count++]->instr);
2882 }
2883 }
2884 }
2885
2886 return str;
2887 }
2888