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