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 fprintf(fp, "call %s ", instr->callee->name);
1909
1910 for (unsigned i = 0; i < instr->num_params; i++) {
1911 if (i != 0)
1912 fprintf(fp, ", ");
1913
1914 if (instr->callee->params[i].name)
1915 fprintf(fp, "%s ", instr->callee->params[i].name);
1916
1917 print_src(&instr->params[i], state, nir_type_invalid);
1918 }
1919 }
1920
1921 static void
print_jump_instr(nir_jump_instr * instr,print_state * state)1922 print_jump_instr(nir_jump_instr *instr, print_state *state)
1923 {
1924 FILE *fp = state->fp;
1925
1926 print_no_dest_padding(state);
1927
1928 switch (instr->type) {
1929 case nir_jump_break:
1930 fprintf(fp, "break");
1931 break;
1932
1933 case nir_jump_continue:
1934 fprintf(fp, "continue");
1935 break;
1936
1937 case nir_jump_return:
1938 fprintf(fp, "return");
1939 break;
1940
1941 case nir_jump_halt:
1942 fprintf(fp, "halt");
1943 break;
1944
1945 case nir_jump_goto:
1946 fprintf(fp, "goto b%u",
1947 instr->target ? instr->target->index : -1);
1948 break;
1949
1950 case nir_jump_goto_if:
1951 fprintf(fp, "goto b%u if ",
1952 instr->target ? instr->target->index : -1);
1953 print_src(&instr->condition, state, nir_type_invalid);
1954 fprintf(fp, " else b%u",
1955 instr->else_target ? instr->else_target->index : -1);
1956 break;
1957 }
1958 }
1959
1960 static void
print_ssa_undef_instr(nir_undef_instr * instr,print_state * state)1961 print_ssa_undef_instr(nir_undef_instr *instr, print_state *state)
1962 {
1963 FILE *fp = state->fp;
1964 print_def(&instr->def, state);
1965 fprintf(fp, " = undefined");
1966 }
1967
1968 static void
print_phi_instr(nir_phi_instr * instr,print_state * state)1969 print_phi_instr(nir_phi_instr *instr, print_state *state)
1970 {
1971 FILE *fp = state->fp;
1972 print_def(&instr->def, state);
1973 fprintf(fp, " = phi ");
1974 nir_foreach_phi_src(src, instr) {
1975 if (&src->node != exec_list_get_head(&instr->srcs))
1976 fprintf(fp, ", ");
1977
1978 fprintf(fp, "b%u: ", src->pred->index);
1979 print_src(&src->src, state, nir_type_invalid);
1980 }
1981 }
1982
1983 static void
print_parallel_copy_instr(nir_parallel_copy_instr * instr,print_state * state)1984 print_parallel_copy_instr(nir_parallel_copy_instr *instr, print_state *state)
1985 {
1986 FILE *fp = state->fp;
1987 nir_foreach_parallel_copy_entry(entry, instr) {
1988 if (&entry->node != exec_list_get_head(&instr->entries))
1989 fprintf(fp, "; ");
1990
1991 if (entry->dest_is_reg) {
1992 fprintf(fp, "*");
1993 print_src(&entry->dest.reg, state, nir_type_invalid);
1994 } else {
1995 print_def(&entry->dest.def, state);
1996 }
1997 fprintf(fp, " = ");
1998
1999 if (entry->src_is_reg)
2000 fprintf(fp, "*");
2001 print_src(&entry->src, state, nir_type_invalid);
2002 }
2003 }
2004
2005 static void
print_debug_info_instr(nir_debug_info_instr * instr,print_state * state)2006 print_debug_info_instr(nir_debug_info_instr *instr, print_state *state)
2007 {
2008 FILE *fp = state->fp;
2009
2010 switch (instr->type) {
2011 case nir_debug_info_src_loc:
2012 fprintf(fp, "// 0x%x", instr->src_loc.spirv_offset);
2013 if (instr->src_loc.line)
2014 fprintf(fp, " %s:%u:%u", nir_src_as_string(instr->src_loc.filename), instr->src_loc.line, instr->src_loc.column);
2015 return;
2016 case nir_debug_info_string:
2017 return; /* Strings are printed for their uses. */
2018 }
2019
2020 unreachable("Unimplemented nir_debug_info_type");
2021 }
2022
2023 static void
print_instr(const nir_instr * instr,print_state * state,unsigned tabs)2024 print_instr(const nir_instr *instr, print_state *state, unsigned tabs)
2025 {
2026 FILE *fp = state->fp;
2027
2028 if (state->debug_info) {
2029 nir_debug_info_instr *di = state->debug_info[instr->index];
2030 if (di)
2031 di->src_loc.column = (uint32_t)ftell(fp);
2032 }
2033
2034 print_indentation(tabs, fp);
2035
2036 switch (instr->type) {
2037 case nir_instr_type_alu:
2038 print_alu_instr(nir_instr_as_alu(instr), state);
2039 break;
2040
2041 case nir_instr_type_deref:
2042 print_deref_instr(nir_instr_as_deref(instr), state);
2043 break;
2044
2045 case nir_instr_type_call:
2046 print_call_instr(nir_instr_as_call(instr), state);
2047 break;
2048
2049 case nir_instr_type_intrinsic:
2050 print_intrinsic_instr(nir_instr_as_intrinsic(instr), state);
2051 break;
2052
2053 case nir_instr_type_tex:
2054 print_tex_instr(nir_instr_as_tex(instr), state);
2055 break;
2056
2057 case nir_instr_type_load_const:
2058 print_load_const_instr(nir_instr_as_load_const(instr), state);
2059 break;
2060
2061 case nir_instr_type_jump:
2062 print_jump_instr(nir_instr_as_jump(instr), state);
2063 break;
2064
2065 case nir_instr_type_undef:
2066 print_ssa_undef_instr(nir_instr_as_undef(instr), state);
2067 break;
2068
2069 case nir_instr_type_phi:
2070 print_phi_instr(nir_instr_as_phi(instr), state);
2071 break;
2072
2073 case nir_instr_type_parallel_copy:
2074 print_parallel_copy_instr(nir_instr_as_parallel_copy(instr), state);
2075 break;
2076
2077 case nir_instr_type_debug_info:
2078 print_debug_info_instr(nir_instr_as_debug_info(instr), state);
2079 break;
2080
2081 default:
2082 unreachable("Invalid instruction type");
2083 break;
2084 }
2085
2086 if (NIR_DEBUG(PRINT_PASS_FLAGS) && instr->pass_flags)
2087 fprintf(fp, " (pass_flags: 0x%x)", instr->pass_flags);
2088 }
2089
2090 static bool
block_has_instruction_with_dest(nir_block * block)2091 block_has_instruction_with_dest(nir_block *block)
2092 {
2093 nir_foreach_instr(instr, block) {
2094 switch (instr->type) {
2095 case nir_instr_type_load_const:
2096 case nir_instr_type_deref:
2097 case nir_instr_type_alu:
2098 case nir_instr_type_tex:
2099 case nir_instr_type_undef:
2100 case nir_instr_type_phi:
2101 case nir_instr_type_parallel_copy:
2102 return true;
2103
2104 case nir_instr_type_intrinsic: {
2105 nir_intrinsic_instr *intrin = nir_instr_as_intrinsic(instr);
2106 const nir_intrinsic_info *info = &nir_intrinsic_infos[intrin->intrinsic];
2107 if (info->has_dest)
2108 return true;
2109
2110 /* Doesn't define a new value. */
2111 break;
2112 }
2113
2114 case nir_instr_type_jump:
2115 case nir_instr_type_call:
2116 case nir_instr_type_debug_info:
2117 /* Doesn't define a new value. */
2118 break;
2119 }
2120 }
2121
2122 return false;
2123 }
2124
2125 static void print_cf_node(nir_cf_node *node, print_state *state,
2126 unsigned tabs);
2127
2128 static void
print_block_preds(nir_block * block,print_state * state)2129 print_block_preds(nir_block *block, print_state *state)
2130 {
2131 FILE *fp = state->fp;
2132 nir_block **preds = nir_block_get_predecessors_sorted(block, NULL);
2133 for (unsigned i = 0; i < block->predecessors->entries; i++) {
2134 if (i != 0)
2135 fprintf(fp, " ");
2136 fprintf(fp, "b%u", preds[i]->index);
2137 }
2138 ralloc_free(preds);
2139 }
2140
2141 static void
print_block_succs(nir_block * block,print_state * state)2142 print_block_succs(nir_block *block, print_state *state)
2143 {
2144 FILE *fp = state->fp;
2145 for (unsigned i = 0; i < 2; i++) {
2146 if (block->successors[i]) {
2147 fprintf(fp, "b%u ", block->successors[i]->index);
2148 }
2149 }
2150 }
2151
2152 static void
print_block(nir_block * block,print_state * state,unsigned tabs)2153 print_block(nir_block *block, print_state *state, unsigned tabs)
2154 {
2155 FILE *fp = state->fp;
2156
2157 if (block_has_instruction_with_dest(block))
2158 state->padding_for_no_dest = calculate_padding_for_no_dest(state);
2159 else
2160 state->padding_for_no_dest = 0;
2161
2162 print_indentation(tabs, fp);
2163 fprintf(fp, "%sblock b%u:",
2164 divergence_status(state, block->divergent),
2165 block->index);
2166
2167 const bool empty_block = exec_list_is_empty(&block->instr_list);
2168 if (empty_block) {
2169 fprintf(fp, " // preds: ");
2170 print_block_preds(block, state);
2171 fprintf(fp, ", succs: ");
2172 print_block_succs(block, state);
2173 fprintf(fp, "\n");
2174 return;
2175 }
2176
2177 const unsigned block_length = 7 + count_digits(block->index) + 1;
2178 const unsigned pred_padding = block_length < state->padding_for_no_dest ? state->padding_for_no_dest - block_length : 0;
2179
2180 fprintf(fp, "%*s// preds: ", pred_padding, "");
2181 print_block_preds(block, state);
2182 fprintf(fp, "\n");
2183
2184 nir_foreach_instr(instr, block) {
2185 print_instr(instr, state, tabs);
2186 fprintf(fp, "\n");
2187 print_annotation(state, instr);
2188 }
2189
2190 print_indentation(tabs, fp);
2191 fprintf(fp, "%*s// succs: ", state->padding_for_no_dest, "");
2192 print_block_succs(block, state);
2193 fprintf(fp, "\n");
2194 }
2195
2196 static void
print_if(nir_if * if_stmt,print_state * state,unsigned tabs)2197 print_if(nir_if *if_stmt, print_state *state, unsigned tabs)
2198 {
2199 FILE *fp = state->fp;
2200
2201 print_indentation(tabs, fp);
2202 fprintf(fp, "if ");
2203 print_src(&if_stmt->condition, state, nir_type_invalid);
2204 switch (if_stmt->control) {
2205 case nir_selection_control_flatten:
2206 fprintf(fp, " // flatten");
2207 break;
2208 case nir_selection_control_dont_flatten:
2209 fprintf(fp, " // don't flatten");
2210 break;
2211 case nir_selection_control_divergent_always_taken:
2212 fprintf(fp, " // divergent always taken");
2213 break;
2214 case nir_selection_control_none:
2215 default:
2216 break;
2217 }
2218 fprintf(fp, " {\n");
2219 foreach_list_typed(nir_cf_node, node, node, &if_stmt->then_list) {
2220 print_cf_node(node, state, tabs + 1);
2221 }
2222 print_indentation(tabs, fp);
2223 fprintf(fp, "} else {\n");
2224 foreach_list_typed(nir_cf_node, node, node, &if_stmt->else_list) {
2225 print_cf_node(node, state, tabs + 1);
2226 }
2227 print_indentation(tabs, fp);
2228 fprintf(fp, "}\n");
2229 }
2230
2231 static void
print_loop(nir_loop * loop,print_state * state,unsigned tabs)2232 print_loop(nir_loop *loop, print_state *state, unsigned tabs)
2233 {
2234 FILE *fp = state->fp;
2235
2236 print_indentation(tabs, fp);
2237 fprintf(fp, "%sloop {\n", divergence_status(state, loop->divergent_break));
2238 foreach_list_typed(nir_cf_node, node, node, &loop->body) {
2239 print_cf_node(node, state, tabs + 1);
2240 }
2241 print_indentation(tabs, fp);
2242
2243 if (nir_loop_has_continue_construct(loop)) {
2244 fprintf(fp, "} continue {\n");
2245 foreach_list_typed(nir_cf_node, node, node, &loop->continue_list) {
2246 print_cf_node(node, state, tabs + 1);
2247 }
2248 print_indentation(tabs, fp);
2249 }
2250
2251 fprintf(fp, "}\n");
2252 }
2253
2254 static void
print_cf_node(nir_cf_node * node,print_state * state,unsigned int tabs)2255 print_cf_node(nir_cf_node *node, print_state *state, unsigned int tabs)
2256 {
2257 switch (node->type) {
2258 case nir_cf_node_block:
2259 print_block(nir_cf_node_as_block(node), state, tabs);
2260 break;
2261
2262 case nir_cf_node_if:
2263 print_if(nir_cf_node_as_if(node), state, tabs);
2264 break;
2265
2266 case nir_cf_node_loop:
2267 print_loop(nir_cf_node_as_loop(node), state, tabs);
2268 break;
2269
2270 default:
2271 unreachable("Invalid CFG node type");
2272 }
2273 }
2274
2275 static void
print_function_impl(nir_function_impl * impl,print_state * state)2276 print_function_impl(nir_function_impl *impl, print_state *state)
2277 {
2278 FILE *fp = state->fp;
2279
2280 state->max_dest_index = impl->ssa_alloc;
2281
2282 fprintf(fp, "\nimpl %s ", impl->function->name);
2283
2284 fprintf(fp, "{\n");
2285
2286 if (impl->preamble) {
2287 print_indentation(1, fp);
2288 fprintf(fp, "preamble %s\n", impl->preamble->name);
2289 }
2290
2291 if (!NIR_DEBUG(PRINT_NO_INLINE_CONSTS)) {
2292 /* Don't reindex the SSA as suggested by nir_gather_types() because
2293 * nir_print don't modify the shader. If needed, a limit for ssa_alloc
2294 * can be added.
2295 */
2296 state->float_types = calloc(BITSET_WORDS(impl->ssa_alloc), sizeof(BITSET_WORD));
2297 state->int_types = calloc(BITSET_WORDS(impl->ssa_alloc), sizeof(BITSET_WORD));
2298 nir_gather_types(impl, state->float_types, state->int_types);
2299 }
2300
2301 nir_foreach_function_temp_variable(var, impl) {
2302 print_indentation(1, fp);
2303 print_var_decl(var, state);
2304 }
2305
2306 nir_index_blocks(impl);
2307
2308 foreach_list_typed(nir_cf_node, node, node, &impl->body) {
2309 print_cf_node(node, state, 1);
2310 }
2311
2312 print_indentation(1, fp);
2313 fprintf(fp, "block b%u:\n}\n\n", impl->end_block->index);
2314
2315 free(state->float_types);
2316 free(state->int_types);
2317 state->max_dest_index = 0;
2318 }
2319
2320 static void
print_function(nir_function * function,print_state * state)2321 print_function(nir_function *function, print_state *state)
2322 {
2323 FILE *fp = state->fp;
2324
2325 fprintf(fp, "decl_function %s (", function->name);
2326
2327 for (unsigned i = 0; i < function->num_params; ++i) {
2328 if (i != 0) {
2329 fprintf(fp, ", ");
2330 }
2331
2332 nir_parameter param = function->params[i];
2333
2334 fprintf(fp, "%u", param.bit_size);
2335 if (param.num_components != 1) {
2336 fprintf(fp, "x%u", param.num_components);
2337 }
2338
2339 if (param.name) {
2340 fprintf(fp, " %s", param.name);
2341 } else if (param.is_return) {
2342 fprintf(fp, " return");
2343 }
2344 }
2345
2346 fprintf(fp, ")");
2347
2348 /* clang-format off */
2349 fprintf(fp, "%s%s%s", function->dont_inline ? " (noinline)" :
2350 function->should_inline ? " (inline)" : "",
2351 function->is_exported ? " (exported)" : "",
2352 function->is_entrypoint ? " (entrypoint)" : "");
2353 /* clang-format on */
2354
2355 if (function->workgroup_size[0]) {
2356 fprintf(fp, " (%ux%ux%u)",
2357 function->workgroup_size[0],
2358 function->workgroup_size[1],
2359 function->workgroup_size[2]);
2360 }
2361
2362 fprintf(fp, "\n");
2363
2364 if (function->impl != NULL) {
2365 print_function_impl(function->impl, state);
2366 return;
2367 }
2368 }
2369
2370 static void
init_print_state(print_state * state,nir_shader * shader,FILE * fp)2371 init_print_state(print_state *state, nir_shader *shader, FILE *fp)
2372 {
2373 state->fp = fp;
2374 state->shader = shader;
2375 state->ht = _mesa_pointer_hash_table_create(NULL);
2376 state->syms = _mesa_set_create(NULL, _mesa_hash_string,
2377 _mesa_key_string_equal);
2378 state->index = 0;
2379 state->int_types = NULL;
2380 state->float_types = NULL;
2381 state->max_dest_index = 0;
2382 state->padding_for_no_dest = 0;
2383 }
2384
2385 static void
destroy_print_state(print_state * state)2386 destroy_print_state(print_state *state)
2387 {
2388 _mesa_hash_table_destroy(state->ht, NULL);
2389 _mesa_set_destroy(state->syms, NULL);
2390 }
2391
2392 static const char *
primitive_name(unsigned primitive)2393 primitive_name(unsigned primitive)
2394 {
2395 #define PRIM(X) \
2396 case MESA_PRIM_##X: \
2397 return #X
2398 switch (primitive) {
2399 PRIM(POINTS);
2400 PRIM(LINES);
2401 PRIM(LINE_LOOP);
2402 PRIM(LINE_STRIP);
2403 PRIM(TRIANGLES);
2404 PRIM(TRIANGLE_STRIP);
2405 PRIM(TRIANGLE_FAN);
2406 PRIM(QUADS);
2407 PRIM(QUAD_STRIP);
2408 PRIM(POLYGON);
2409 PRIM(LINES_ADJACENCY);
2410 PRIM(TRIANGLES_ADJACENCY);
2411 default:
2412 return "UNKNOWN";
2413 }
2414 }
2415
2416 static void
print_bitset(FILE * fp,const char * label,const unsigned * words,int size)2417 print_bitset(FILE *fp, const char *label, const unsigned *words, int size)
2418 {
2419 fprintf(fp, "%s: ", label);
2420 /* Iterate back-to-front to get proper digit order (most significant first). */
2421 for (int i = size - 1; i >= 0; --i) {
2422 fprintf(fp, (i == size - 1) ? "0x%08x" : "'%08x", words[i]);
2423 }
2424 fprintf(fp, "\n");
2425 }
2426
2427 /* Print bitset, only if some bits are set */
2428 static void
print_nz_bitset(FILE * fp,const char * label,const unsigned * words,int size)2429 print_nz_bitset(FILE *fp, const char *label, const unsigned *words, int size)
2430 {
2431 bool is_all_zero = true;
2432 for (int i = 0; i < size; ++i) {
2433 if (words[i]) {
2434 is_all_zero = false;
2435 break;
2436 }
2437 }
2438
2439 if (!is_all_zero)
2440 print_bitset(fp, label, words, size);
2441 }
2442
2443 /* Print uint64_t value, only if non-zero.
2444 * The value is printed by enumerating the ranges of bits that are set.
2445 * E.g. inputs_read: 0,15-17
2446 */
2447 static void
print_nz_x64(FILE * fp,const char * label,uint64_t value)2448 print_nz_x64(FILE *fp, const char *label, uint64_t value)
2449 {
2450 if (value) {
2451 char acc[256] = { 0 };
2452 char buf[32];
2453 int start = 0;
2454 int count = 0;
2455 while (value) {
2456 u_bit_scan_consecutive_range64(&value, &start, &count);
2457 assert(count > 0);
2458 bool is_first = !acc[0];
2459 if (count > 1) {
2460 snprintf(buf, sizeof(buf), is_first ? "%d-%d" : ",%d-%d", start, start + count - 1);
2461 } else {
2462 snprintf(buf, sizeof(buf), is_first ? "%d" : ",%d", start);
2463 }
2464 assert(strlen(acc) + strlen(buf) + 1 < sizeof(acc));
2465 strcat(acc, buf);
2466 }
2467 fprintf(fp, "%s: %s\n", label, acc);
2468 }
2469 }
2470
2471 /* Print uint32_t value in hex, only if non-zero */
2472 static void
print_nz_x32(FILE * fp,const char * label,uint32_t value)2473 print_nz_x32(FILE *fp, const char *label, uint32_t value)
2474 {
2475 if (value)
2476 fprintf(fp, "%s: 0x%08" PRIx32 "\n", label, value);
2477 }
2478
2479 /* Print uint16_t value in hex, only if non-zero */
2480 static void
print_nz_x16(FILE * fp,const char * label,uint16_t value)2481 print_nz_x16(FILE *fp, const char *label, uint16_t value)
2482 {
2483 if (value)
2484 fprintf(fp, "%s: 0x%04x\n", label, value);
2485 }
2486
2487 /* Print uint8_t value in hex, only if non-zero */
2488 static void
print_nz_x8(FILE * fp,const char * label,uint8_t value)2489 print_nz_x8(FILE *fp, const char *label, uint8_t value)
2490 {
2491 if (value)
2492 fprintf(fp, "%s: 0x%02x\n", label, value);
2493 }
2494
2495 /* Print unsigned value in decimal, only if non-zero */
2496 static void
print_nz_unsigned(FILE * fp,const char * label,unsigned value)2497 print_nz_unsigned(FILE *fp, const char *label, unsigned value)
2498 {
2499 if (value)
2500 fprintf(fp, "%s: %u\n", label, value);
2501 }
2502
2503 /* Print bool only if set */
2504 static void
print_nz_bool(FILE * fp,const char * label,bool value)2505 print_nz_bool(FILE *fp, const char *label, bool value)
2506 {
2507 if (value)
2508 fprintf(fp, "%s: true\n", label);
2509 }
2510
2511 static void
print_shader_info(const struct shader_info * info,FILE * fp)2512 print_shader_info(const struct shader_info *info, FILE *fp)
2513 {
2514 fprintf(fp, "shader: %s\n", gl_shader_stage_name(info->stage));
2515
2516 fprintf(fp, "source_blake3: {");
2517 _mesa_blake3_print(fp, info->source_blake3);
2518 fprintf(fp, "}\n");
2519
2520 if (info->name)
2521 fprintf(fp, "name: %s\n", info->name);
2522
2523 if (info->label)
2524 fprintf(fp, "label: %s\n", info->label);
2525
2526 fprintf(fp, "internal: %s\n", info->internal ? "true" : "false");
2527
2528 if (gl_shader_stage_uses_workgroup(info->stage)) {
2529 fprintf(fp, "workgroup_size: %u, %u, %u%s\n",
2530 info->workgroup_size[0],
2531 info->workgroup_size[1],
2532 info->workgroup_size[2],
2533 info->workgroup_size_variable ? " (variable)" : "");
2534 }
2535
2536 fprintf(fp, "stage: %d\n"
2537 "next_stage: %d\n",
2538 info->stage, info->next_stage);
2539
2540 print_nz_unsigned(fp, "num_textures", info->num_textures);
2541 print_nz_unsigned(fp, "num_ubos", info->num_ubos);
2542 print_nz_unsigned(fp, "num_abos", info->num_abos);
2543 print_nz_unsigned(fp, "num_ssbos", info->num_ssbos);
2544 print_nz_unsigned(fp, "num_images", info->num_images);
2545
2546 print_nz_x64(fp, "inputs_read", info->inputs_read);
2547 print_nz_x64(fp, "dual_slot_inputs", info->dual_slot_inputs);
2548 print_nz_x64(fp, "outputs_written", info->outputs_written);
2549 print_nz_x64(fp, "outputs_read", info->outputs_read);
2550
2551 print_nz_bitset(fp, "system_values_read", info->system_values_read, ARRAY_SIZE(info->system_values_read));
2552
2553 print_nz_x64(fp, "per_primitive_inputs", info->per_primitive_inputs);
2554 print_nz_x64(fp, "per_primitive_outputs", info->per_primitive_outputs);
2555 print_nz_x64(fp, "per_view_outputs", info->per_view_outputs);
2556
2557 print_nz_x16(fp, "inputs_read_16bit", info->inputs_read_16bit);
2558 print_nz_x16(fp, "outputs_written_16bit", info->outputs_written_16bit);
2559 print_nz_x16(fp, "outputs_read_16bit", info->outputs_read_16bit);
2560 print_nz_x16(fp, "inputs_read_indirectly_16bit", info->inputs_read_indirectly_16bit);
2561 print_nz_x16(fp, "outputs_accessed_indirectly_16bit", info->outputs_accessed_indirectly_16bit);
2562
2563 print_nz_x32(fp, "patch_inputs_read", info->patch_inputs_read);
2564 print_nz_x32(fp, "patch_outputs_written", info->patch_outputs_written);
2565 print_nz_x32(fp, "patch_outputs_read", info->patch_outputs_read);
2566
2567 print_nz_x64(fp, "inputs_read_indirectly", info->inputs_read_indirectly);
2568 print_nz_x64(fp, "outputs_accessed_indirectly", info->outputs_accessed_indirectly);
2569 print_nz_x64(fp, "patch_inputs_read_indirectly", info->patch_inputs_read_indirectly);
2570 print_nz_x64(fp, "patch_outputs_accessed_indirectly", info->patch_outputs_accessed_indirectly);
2571
2572 print_nz_bitset(fp, "textures_used", info->textures_used, ARRAY_SIZE(info->textures_used));
2573 print_nz_bitset(fp, "textures_used_by_txf", info->textures_used_by_txf, ARRAY_SIZE(info->textures_used_by_txf));
2574 print_nz_bitset(fp, "samplers_used", info->samplers_used, ARRAY_SIZE(info->samplers_used));
2575 print_nz_bitset(fp, "images_used", info->images_used, ARRAY_SIZE(info->images_used));
2576 print_nz_bitset(fp, "image_buffers", info->image_buffers, ARRAY_SIZE(info->image_buffers));
2577 print_nz_bitset(fp, "msaa_images", info->msaa_images, ARRAY_SIZE(info->msaa_images));
2578
2579 print_nz_x32(fp, "float_controls_execution_mode", info->float_controls_execution_mode);
2580
2581 print_nz_unsigned(fp, "shared_size", info->shared_size);
2582
2583 if (info->stage == MESA_SHADER_MESH || info->stage == MESA_SHADER_TASK) {
2584 fprintf(fp, "task_payload_size: %u\n", info->task_payload_size);
2585 }
2586
2587 print_nz_unsigned(fp, "ray queries", info->ray_queries);
2588
2589 fprintf(fp, "subgroup_size: %u\n", info->subgroup_size);
2590
2591 print_nz_bool(fp, "uses_wide_subgroup_intrinsics", info->uses_wide_subgroup_intrinsics);
2592
2593 bool has_xfb_stride = info->xfb_stride[0] || info->xfb_stride[1] || info->xfb_stride[2] || info->xfb_stride[3];
2594 if (has_xfb_stride)
2595 fprintf(fp, "xfb_stride: {%u, %u, %u, %u}\n",
2596 info->xfb_stride[0],
2597 info->xfb_stride[1],
2598 info->xfb_stride[2],
2599 info->xfb_stride[3]);
2600
2601 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];
2602 if (has_inlinable_uniform_dw_offsets)
2603 fprintf(fp, "inlinable_uniform_dw_offsets: {%u, %u, %u, %u}\n",
2604 info->inlinable_uniform_dw_offsets[0],
2605 info->inlinable_uniform_dw_offsets[1],
2606 info->inlinable_uniform_dw_offsets[2],
2607 info->inlinable_uniform_dw_offsets[3]);
2608
2609 print_nz_unsigned(fp, "num_inlinable_uniforms", info->num_inlinable_uniforms);
2610 print_nz_unsigned(fp, "clip_distance_array_size", info->clip_distance_array_size);
2611 print_nz_unsigned(fp, "cull_distance_array_size", info->cull_distance_array_size);
2612
2613 print_nz_bool(fp, "uses_texture_gather", info->uses_texture_gather);
2614 print_nz_bool(fp, "uses_resource_info_query", info->uses_resource_info_query);
2615 print_nz_bool(fp, "divergence_analysis_run", info->divergence_analysis_run);
2616
2617 print_nz_x8(fp, "bit_sizes_float", info->bit_sizes_float);
2618 print_nz_x8(fp, "bit_sizes_int", info->bit_sizes_int);
2619
2620 print_nz_bool(fp, "first_ubo_is_default_ubo", info->first_ubo_is_default_ubo);
2621 print_nz_bool(fp, "separate_shader", info->separate_shader);
2622 print_nz_bool(fp, "has_transform_feedback_varyings", info->has_transform_feedback_varyings);
2623 print_nz_bool(fp, "flrp_lowered", info->flrp_lowered);
2624 print_nz_bool(fp, "io_lowered", info->io_lowered);
2625 print_nz_bool(fp, "writes_memory", info->writes_memory);
2626 print_nz_unsigned(fp, "derivative_group", info->derivative_group);
2627
2628 switch (info->stage) {
2629 case MESA_SHADER_VERTEX:
2630 print_nz_x64(fp, "double_inputs", info->vs.double_inputs);
2631 print_nz_unsigned(fp, "blit_sgprs_amd", info->vs.blit_sgprs_amd);
2632 print_nz_bool(fp, "window_space_position", info->vs.window_space_position);
2633 print_nz_bool(fp, "needs_edge_flag", info->vs.needs_edge_flag);
2634 break;
2635
2636 case MESA_SHADER_TESS_CTRL:
2637 case MESA_SHADER_TESS_EVAL:
2638 fprintf(fp, "primitive_mode: %u\n", info->tess._primitive_mode);
2639 fprintf(fp, "tcs_vertices_out: %u\n", info->tess.tcs_vertices_out);
2640 fprintf(fp, "spacing: %u\n", info->tess.spacing);
2641
2642 print_nz_bool(fp, "ccw", info->tess.ccw);
2643 print_nz_bool(fp, "point_mode", info->tess.point_mode);
2644 print_nz_x64(fp, "tcs_same_invocation_inputs_read",
2645 info->tess.tcs_same_invocation_inputs_read);
2646 print_nz_x64(fp, "tcs_cross_invocation_inputs_read", info->tess.tcs_cross_invocation_inputs_read);
2647 print_nz_x64(fp, "tcs_cross_invocation_outputs_read", info->tess.tcs_cross_invocation_outputs_read);
2648 break;
2649
2650 case MESA_SHADER_GEOMETRY:
2651 fprintf(fp, "output_primitive: %s\n", primitive_name(info->gs.output_primitive));
2652 fprintf(fp, "input_primitive: %s\n", primitive_name(info->gs.input_primitive));
2653 fprintf(fp, "vertices_out: %u\n", info->gs.vertices_out);
2654 fprintf(fp, "invocations: %u\n", info->gs.invocations);
2655 fprintf(fp, "vertices_in: %u\n", info->gs.vertices_in);
2656 print_nz_bool(fp, "uses_end_primitive", info->gs.uses_end_primitive);
2657 fprintf(fp, "active_stream_mask: 0x%02x\n", info->gs.active_stream_mask);
2658 break;
2659
2660 case MESA_SHADER_FRAGMENT:
2661 print_nz_bool(fp, "uses_discard", info->fs.uses_discard);
2662 print_nz_bool(fp, "uses_fbfetch_output", info->fs.uses_fbfetch_output);
2663 print_nz_bool(fp, "color_is_dual_source", info->fs.color_is_dual_source);
2664
2665 print_nz_bool(fp, "require_full_quads", info->fs.require_full_quads);
2666 print_nz_bool(fp, "needs_quad_helper_invocations", info->fs.needs_quad_helper_invocations);
2667 print_nz_bool(fp, "uses_sample_qualifier", info->fs.uses_sample_qualifier);
2668 print_nz_bool(fp, "uses_sample_shading", info->fs.uses_sample_shading);
2669 print_nz_bool(fp, "early_fragment_tests", info->fs.early_fragment_tests);
2670 print_nz_bool(fp, "inner_coverage", info->fs.inner_coverage);
2671 print_nz_bool(fp, "post_depth_coverage", info->fs.post_depth_coverage);
2672
2673 print_nz_bool(fp, "pixel_center_integer", info->fs.pixel_center_integer);
2674 print_nz_bool(fp, "origin_upper_left", info->fs.origin_upper_left);
2675 print_nz_bool(fp, "pixel_interlock_ordered", info->fs.pixel_interlock_ordered);
2676 print_nz_bool(fp, "pixel_interlock_unordered", info->fs.pixel_interlock_unordered);
2677 print_nz_bool(fp, "sample_interlock_ordered", info->fs.sample_interlock_ordered);
2678 print_nz_bool(fp, "sample_interlock_unordered", info->fs.sample_interlock_unordered);
2679 print_nz_bool(fp, "untyped_color_outputs", info->fs.untyped_color_outputs);
2680
2681 print_nz_unsigned(fp, "depth_layout", info->fs.depth_layout);
2682
2683 if (info->fs.color0_interp != INTERP_MODE_NONE) {
2684 fprintf(fp, "color0_interp: %s\n",
2685 glsl_interp_mode_name(info->fs.color0_interp));
2686 }
2687 print_nz_bool(fp, "color0_sample", info->fs.color0_sample);
2688 print_nz_bool(fp, "color0_centroid", info->fs.color0_centroid);
2689
2690 if (info->fs.color1_interp != INTERP_MODE_NONE) {
2691 fprintf(fp, "color1_interp: %s\n",
2692 glsl_interp_mode_name(info->fs.color1_interp));
2693 }
2694 print_nz_bool(fp, "color1_sample", info->fs.color1_sample);
2695 print_nz_bool(fp, "color1_centroid", info->fs.color1_centroid);
2696
2697 print_nz_x32(fp, "advanced_blend_modes", info->fs.advanced_blend_modes);
2698 break;
2699
2700 case MESA_SHADER_COMPUTE:
2701 case MESA_SHADER_KERNEL:
2702 if (info->cs.workgroup_size_hint[0] || info->cs.workgroup_size_hint[1] || info->cs.workgroup_size_hint[2])
2703 fprintf(fp, "workgroup_size_hint: {%u, %u, %u}\n",
2704 info->cs.workgroup_size_hint[0],
2705 info->cs.workgroup_size_hint[1],
2706 info->cs.workgroup_size_hint[2]);
2707 print_nz_unsigned(fp, "user_data_components_amd", info->cs.user_data_components_amd);
2708 fprintf(fp, "ptr_size: %u\n", info->cs.ptr_size);
2709 break;
2710
2711 case MESA_SHADER_MESH:
2712 print_nz_x64(fp, "ms_cross_invocation_output_access", info->mesh.ms_cross_invocation_output_access);
2713 fprintf(fp, "max_vertices_out: %u\n", info->mesh.max_vertices_out);
2714 fprintf(fp, "max_primitives_out: %u\n", info->mesh.max_primitives_out);
2715 fprintf(fp, "primitive_type: %s\n", primitive_name(info->mesh.primitive_type));
2716 print_nz_bool(fp, "nv", info->mesh.nv);
2717 break;
2718
2719 default:
2720 fprintf(fp, "Unhandled stage %d\n", info->stage);
2721 }
2722 }
2723
2724 static void
_nir_print_shader_annotated(nir_shader * shader,FILE * fp,struct hash_table * annotations,nir_debug_info_instr ** debug_info)2725 _nir_print_shader_annotated(nir_shader *shader, FILE *fp,
2726 struct hash_table *annotations,
2727 nir_debug_info_instr **debug_info)
2728 {
2729 print_state state;
2730 init_print_state(&state, shader, fp);
2731 state.def_prefix = debug_info ? "ssa_" : "%";
2732 state.annotations = annotations;
2733 state.debug_info = debug_info;
2734
2735 print_shader_info(&shader->info, fp);
2736
2737 fprintf(fp, "inputs: %u\n", shader->num_inputs);
2738 fprintf(fp, "outputs: %u\n", shader->num_outputs);
2739 fprintf(fp, "uniforms: %u\n", shader->num_uniforms);
2740 if (shader->scratch_size)
2741 fprintf(fp, "scratch: %u\n", shader->scratch_size);
2742 if (shader->constant_data_size)
2743 fprintf(fp, "constants: %u\n", shader->constant_data_size);
2744 for (unsigned i = 0; i < nir_num_variable_modes; i++) {
2745 nir_variable_mode mode = BITFIELD_BIT(i);
2746 if (mode == nir_var_function_temp)
2747 continue;
2748
2749 if (mode == nir_var_shader_in || mode == nir_var_shader_out) {
2750 for (unsigned j = 0; j < 128; j++) {
2751 nir_variable *vars[NIR_MAX_VEC_COMPONENTS] = {0};
2752 nir_foreach_variable_with_modes(var, shader, mode) {
2753 if (var->data.location == j)
2754 vars[var->data.location_frac] = var;
2755 }
2756 for (unsigned j = 0; j < ARRAY_SIZE(vars); j++)
2757 if (vars[j]) {
2758 print_var_decl(vars[j], &state);
2759 }
2760 }
2761 } else {
2762 nir_foreach_variable_with_modes(var, shader, mode)
2763 print_var_decl(var, &state);
2764 }
2765 }
2766
2767 foreach_list_typed(nir_function, func, node, &shader->functions) {
2768 print_function(func, &state);
2769 }
2770
2771 destroy_print_state(&state);
2772 }
2773
2774 void
nir_print_shader_annotated(nir_shader * shader,FILE * fp,struct hash_table * annotations)2775 nir_print_shader_annotated(nir_shader *shader, FILE *fp,
2776 struct hash_table *annotations)
2777 {
2778 _nir_print_shader_annotated(shader, fp, annotations, NULL);
2779 }
2780
2781 void
nir_print_shader(nir_shader * shader,FILE * fp)2782 nir_print_shader(nir_shader *shader, FILE *fp)
2783 {
2784 nir_print_shader_annotated(shader, fp, NULL);
2785 fflush(fp);
2786 }
2787
2788 static char *
_nir_shader_as_str_annotated(nir_shader * nir,struct hash_table * annotations,void * mem_ctx,nir_debug_info_instr ** debug_info)2789 _nir_shader_as_str_annotated(nir_shader *nir, struct hash_table *annotations, void *mem_ctx,
2790 nir_debug_info_instr **debug_info)
2791 {
2792 char *stream_data = NULL;
2793 size_t stream_size = 0;
2794 struct u_memstream mem;
2795 if (u_memstream_open(&mem, &stream_data, &stream_size)) {
2796 FILE *const stream = u_memstream_get(&mem);
2797 _nir_print_shader_annotated(nir, stream, annotations, debug_info);
2798 u_memstream_close(&mem);
2799 }
2800
2801 char *str = ralloc_size(mem_ctx, stream_size + 1);
2802 memcpy(str, stream_data, stream_size);
2803 str[stream_size] = '\0';
2804
2805 free(stream_data);
2806
2807 return str;
2808 }
2809
2810 char *
nir_shader_as_str_annotated(nir_shader * nir,struct hash_table * annotations,void * mem_ctx)2811 nir_shader_as_str_annotated(nir_shader *nir, struct hash_table *annotations, void *mem_ctx)
2812 {
2813 return _nir_shader_as_str_annotated(nir, annotations, mem_ctx, NULL);
2814 }
2815
2816 char *
nir_shader_as_str(nir_shader * nir,void * mem_ctx)2817 nir_shader_as_str(nir_shader *nir, void *mem_ctx)
2818 {
2819 return nir_shader_as_str_annotated(nir, NULL, mem_ctx);
2820 }
2821
2822 void
nir_print_instr(const nir_instr * instr,FILE * fp)2823 nir_print_instr(const nir_instr *instr, FILE *fp)
2824 {
2825 print_state state = {
2826 .fp = fp,
2827 .def_prefix = "%",
2828 };
2829 if (instr->block) {
2830 nir_function_impl *impl = nir_cf_node_get_function(&instr->block->cf_node);
2831 state.shader = impl->function->shader;
2832 }
2833
2834 print_instr(instr, &state, 0);
2835 }
2836
2837 char *
nir_instr_as_str(const nir_instr * instr,void * mem_ctx)2838 nir_instr_as_str(const nir_instr *instr, void *mem_ctx)
2839 {
2840 char *stream_data = NULL;
2841 size_t stream_size = 0;
2842 struct u_memstream mem;
2843 if (u_memstream_open(&mem, &stream_data, &stream_size)) {
2844 FILE *const stream = u_memstream_get(&mem);
2845 nir_print_instr(instr, stream);
2846 u_memstream_close(&mem);
2847 }
2848
2849 char *str = ralloc_size(mem_ctx, stream_size + 1);
2850 memcpy(str, stream_data, stream_size);
2851 str[stream_size] = '\0';
2852
2853 free(stream_data);
2854
2855 return str;
2856 }
2857
2858 void
nir_print_deref(const nir_deref_instr * deref,FILE * fp)2859 nir_print_deref(const nir_deref_instr *deref, FILE *fp)
2860 {
2861 print_state state = {
2862 .fp = fp,
2863 .def_prefix = "%",
2864 };
2865 print_deref_link(deref, true, &state);
2866 }
2867
2868 void
nir_log_shader_annotated_tagged(enum mesa_log_level level,const char * tag,nir_shader * shader,struct hash_table * annotations)2869 nir_log_shader_annotated_tagged(enum mesa_log_level level, const char *tag,
2870 nir_shader *shader, struct hash_table *annotations)
2871 {
2872 char *str = nir_shader_as_str_annotated(shader, annotations, NULL);
2873 _mesa_log_multiline(level, tag, str);
2874 ralloc_free(str);
2875 }
2876
2877 char *
nir_shader_gather_debug_info(nir_shader * shader,const char * filename,uint32_t first_line)2878 nir_shader_gather_debug_info(nir_shader *shader, const char *filename, uint32_t first_line)
2879 {
2880 uint32_t instr_count = 0;
2881 nir_foreach_function_impl(impl, shader) {
2882 nir_foreach_block(block, impl) {
2883 nir_foreach_instr(instr, block) {
2884 instr->index = instr_count;
2885 instr_count++;
2886 }
2887 }
2888 }
2889
2890 if (!instr_count)
2891 return nir_shader_as_str(shader, NULL);
2892
2893 nir_debug_info_instr **debug_info = rzalloc_array(shader, nir_debug_info_instr *, instr_count);
2894
2895 instr_count = 0;
2896 nir_foreach_function_impl(impl, shader) {
2897 nir_builder b = nir_builder_at(nir_before_cf_list(&impl->body));
2898 nir_def *filename_def = nir_build_string(&b, filename);
2899
2900 nir_foreach_block(block, impl) {
2901 nir_foreach_instr_safe(instr, block) {
2902 if (instr->type == nir_instr_type_debug_info ||
2903 instr->type == nir_instr_type_phi)
2904 continue;
2905
2906 nir_debug_info_instr *di = nir_debug_info_instr_create(shader, nir_debug_info_src_loc, 0);
2907 di->src_loc.filename = nir_src_for_ssa(filename_def);
2908 di->src_loc.source = nir_debug_info_nir;
2909 debug_info[instr_count++] = di;
2910 }
2911 }
2912 }
2913
2914 char *str = _nir_shader_as_str_annotated(shader, NULL, NULL, debug_info);
2915
2916 uint32_t line = first_line;
2917 uint32_t character_index = 0;
2918
2919 for (uint32_t i = 0; i < instr_count; i++) {
2920 nir_debug_info_instr *di = debug_info[i];
2921 if (!di)
2922 continue;
2923
2924 while (character_index < di->src_loc.column) {
2925 if (str[character_index] == '\n')
2926 line++;
2927 character_index++;
2928 }
2929
2930 di->src_loc.line = line;
2931 di->src_loc.column = 0;
2932 }
2933
2934 instr_count = 0;
2935 nir_foreach_function_impl(impl, shader) {
2936 nir_foreach_block(block, impl) {
2937 nir_foreach_instr_safe(instr, block) {
2938 if (instr->type != nir_instr_type_debug_info &&
2939 instr->type != nir_instr_type_phi)
2940 nir_instr_insert_before(instr, &debug_info[instr_count++]->instr);
2941 }
2942 }
2943 }
2944
2945 return str;
2946 }
2947