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 "nir.h"
29 #include "compiler/shader_enums.h"
30 #include "util/half_float.h"
31 #include "util/memstream.h"
32 #include "util/mesa-sha1.h"
33 #include "vulkan/vulkan_core.h"
34 #include <stdio.h>
35 #include <stdlib.h>
36 #include <inttypes.h> /* for PRIx64 macro */
37
38 static void
print_tabs(unsigned num_tabs,FILE * fp)39 print_tabs(unsigned num_tabs, FILE *fp)
40 {
41 for (unsigned i = 0; i < num_tabs; i++)
42 fprintf(fp, "\t");
43 }
44
45 typedef struct {
46 FILE *fp;
47 nir_shader *shader;
48 /** map from nir_variable -> printable name */
49 struct hash_table *ht;
50
51 /** set of names used so far for nir_variables */
52 struct set *syms;
53
54 /* an index used to make new non-conflicting names */
55 unsigned index;
56
57 /**
58 * Optional table of annotations mapping nir object
59 * (such as instr or var) to message to print.
60 */
61 struct hash_table *annotations;
62 } print_state;
63
64 static void
print_annotation(print_state * state,void * obj)65 print_annotation(print_state *state, void *obj)
66 {
67 FILE *fp = state->fp;
68
69 if (!state->annotations)
70 return;
71
72 struct hash_entry *entry = _mesa_hash_table_search(state->annotations, obj);
73 if (!entry)
74 return;
75
76 const char *note = entry->data;
77 _mesa_hash_table_remove(state->annotations, entry);
78
79 fprintf(fp, "%s\n\n", note);
80 }
81
82 static void
print_register(nir_register * reg,print_state * state)83 print_register(nir_register *reg, print_state *state)
84 {
85 FILE *fp = state->fp;
86 fprintf(fp, "r%u", reg->index);
87 }
88
89 static const char *sizes[] = { "error", "vec1", "vec2", "vec3", "vec4",
90 "vec5", "error", "error", "vec8",
91 "error", "error", "error", "error",
92 "error", "error", "error", "vec16"};
93
94 static void
print_register_decl(nir_register * reg,print_state * state)95 print_register_decl(nir_register *reg, print_state *state)
96 {
97 FILE *fp = state->fp;
98 fprintf(fp, "decl_reg %s %u ", sizes[reg->num_components], reg->bit_size);
99 print_register(reg, state);
100 if (reg->num_array_elems != 0)
101 fprintf(fp, "[%u]", reg->num_array_elems);
102 fprintf(fp, "\n");
103 }
104
105 static void
print_ssa_def(nir_ssa_def * def,print_state * state)106 print_ssa_def(nir_ssa_def *def, print_state *state)
107 {
108 FILE *fp = state->fp;
109
110 const char *divergence = "";
111 if (state->shader->info.divergence_analysis_run)
112 divergence = def->divergent ? "div " : "con ";
113
114 fprintf(fp, "%s %2u %sssa_%u", sizes[def->num_components], def->bit_size,
115 divergence, def->index);
116 }
117
118 static void
print_const_from_load(nir_load_const_instr * instr,print_state * state)119 print_const_from_load(nir_load_const_instr *instr, print_state *state)
120 {
121 FILE *fp = state->fp;
122
123 /*
124 * we don't really know the type of the constant (if it will be used as a
125 * float or an int), so just print the raw constant in hex for fidelity
126 * and then print in float again for readability.
127 */
128
129 fprintf(fp, "(");
130
131 for (unsigned i = 0; i < instr->def.num_components; i++) {
132 if (i != 0)
133 fprintf(fp, ", ");
134
135 switch (instr->def.bit_size) {
136 case 64:
137 fprintf(fp, "0x%016" PRIx64, instr->value[i].u64);
138 break;
139 case 32:
140 fprintf(fp, "0x%08x", instr->value[i].u32);
141 break;
142 case 16:
143 fprintf(fp, "0x%04x", instr->value[i].u16);
144 break;
145 case 8:
146 fprintf(fp, "0x%02x", instr->value[i].u8);
147 break;
148 case 1:
149 fprintf(fp, "%s", instr->value[i].b ? "true" : "false");
150 break;
151 }
152 }
153
154 if (instr->def.bit_size > 8) {
155 if (instr->def.num_components > 1)
156 fprintf(fp, ") = (");
157 else
158 fprintf(fp, " = ");
159
160 for (unsigned i = 0; i < instr->def.num_components; i++) {
161 if (i != 0)
162 fprintf(fp, ", ");
163
164 switch (instr->def.bit_size) {
165 case 64:
166 fprintf(fp, "%f", instr->value[i].f64);
167 break;
168 case 32:
169 fprintf(fp, "%f", instr->value[i].f32);
170 break;
171 case 16:
172 fprintf(fp, "%f", _mesa_half_to_float(instr->value[i].u16));
173 break;
174 default:
175 unreachable("unhandled bit size");
176 }
177 }
178 }
179
180 fprintf(fp, ")");
181 }
182
183 static void
print_load_const_instr(nir_load_const_instr * instr,print_state * state)184 print_load_const_instr(nir_load_const_instr *instr, print_state *state)
185 {
186 FILE *fp = state->fp;
187
188 print_ssa_def(&instr->def, state);
189
190 fprintf(fp, " = load_const ");
191
192 print_const_from_load(instr, state);
193 }
194
195 static void
print_ssa_use(nir_ssa_def * def,print_state * state)196 print_ssa_use(nir_ssa_def *def, print_state *state)
197 {
198 FILE *fp = state->fp;
199 fprintf(fp, "ssa_%u", def->index);
200 nir_instr *instr = def->parent_instr;
201 if (instr->type == nir_instr_type_load_const && NIR_DEBUG(PRINT_CONSTS)) {
202 fprintf(fp, " /*");
203 print_const_from_load(nir_instr_as_load_const(instr), state);
204 fprintf(fp, "*/");
205 }
206 }
207
208 static void print_src(const nir_src *src, print_state *state);
209
210 static void
print_reg_src(const nir_reg_src * src,print_state * state)211 print_reg_src(const nir_reg_src *src, print_state *state)
212 {
213 FILE *fp = state->fp;
214 print_register(src->reg, state);
215 if (src->reg->num_array_elems != 0) {
216 fprintf(fp, "[%u", src->base_offset);
217 if (src->indirect != NULL) {
218 fprintf(fp, " + ");
219 print_src(src->indirect, state);
220 }
221 fprintf(fp, "]");
222 }
223 }
224
225 static void
print_reg_dest(nir_reg_dest * dest,print_state * state)226 print_reg_dest(nir_reg_dest *dest, print_state *state)
227 {
228 FILE *fp = state->fp;
229 print_register(dest->reg, state);
230 if (dest->reg->num_array_elems != 0) {
231 fprintf(fp, "[%u", dest->base_offset);
232 if (dest->indirect != NULL) {
233 fprintf(fp, " + ");
234 print_src(dest->indirect, state);
235 }
236 fprintf(fp, "]");
237 }
238 }
239
240 static void
print_src(const nir_src * src,print_state * state)241 print_src(const nir_src *src, print_state *state)
242 {
243 if (src->is_ssa)
244 print_ssa_use(src->ssa, state);
245 else
246 print_reg_src(&src->reg, state);
247 }
248
249 static void
print_dest(nir_dest * dest,print_state * state)250 print_dest(nir_dest *dest, print_state *state)
251 {
252 if (dest->is_ssa)
253 print_ssa_def(&dest->ssa, state);
254 else
255 print_reg_dest(&dest->reg, state);
256 }
257
258 static const char *
comp_mask_string(unsigned num_components)259 comp_mask_string(unsigned num_components)
260 {
261 return (num_components > 4) ? "abcdefghijklmnop" : "xyzw";
262 }
263
264 static void
print_alu_src(nir_alu_instr * instr,unsigned src,print_state * state)265 print_alu_src(nir_alu_instr *instr, unsigned src, print_state *state)
266 {
267 FILE *fp = state->fp;
268
269 if (instr->src[src].negate)
270 fprintf(fp, "-");
271 if (instr->src[src].abs)
272 fprintf(fp, "abs(");
273
274 print_src(&instr->src[src].src, state);
275
276 bool print_swizzle = false;
277 nir_component_mask_t used_channels = 0;
278
279 for (unsigned i = 0; i < NIR_MAX_VEC_COMPONENTS; i++) {
280 if (!nir_alu_instr_channel_used(instr, src, i))
281 continue;
282
283 used_channels++;
284
285 if (instr->src[src].swizzle[i] != i) {
286 print_swizzle = true;
287 break;
288 }
289 }
290
291 unsigned live_channels = nir_src_num_components(instr->src[src].src);
292
293 if (print_swizzle || used_channels != live_channels) {
294 fprintf(fp, ".");
295 for (unsigned i = 0; i < NIR_MAX_VEC_COMPONENTS; i++) {
296 if (!nir_alu_instr_channel_used(instr, src, i))
297 continue;
298
299 fprintf(fp, "%c", comp_mask_string(live_channels)[instr->src[src].swizzle[i]]);
300 }
301 }
302
303 if (instr->src[src].abs)
304 fprintf(fp, ")");
305 }
306
307 static void
print_alu_dest(nir_alu_dest * dest,print_state * state)308 print_alu_dest(nir_alu_dest *dest, print_state *state)
309 {
310 FILE *fp = state->fp;
311 /* we're going to print the saturate modifier later, after the opcode */
312
313 print_dest(&dest->dest, state);
314
315 if (!dest->dest.is_ssa &&
316 dest->write_mask != (1 << dest->dest.reg.reg->num_components) - 1) {
317 unsigned live_channels = dest->dest.reg.reg->num_components;
318 fprintf(fp, ".");
319 for (unsigned i = 0; i < NIR_MAX_VEC_COMPONENTS; i++)
320 if ((dest->write_mask >> i) & 1)
321 fprintf(fp, "%c", comp_mask_string(live_channels)[i]);
322 }
323 }
324
325 static void
print_alu_instr(nir_alu_instr * instr,print_state * state)326 print_alu_instr(nir_alu_instr *instr, print_state *state)
327 {
328 FILE *fp = state->fp;
329
330 print_alu_dest(&instr->dest, state);
331
332 fprintf(fp, " = %s", nir_op_infos[instr->op].name);
333 if (instr->exact)
334 fprintf(fp, "!");
335 if (instr->dest.saturate)
336 fprintf(fp, ".sat");
337 if (instr->no_signed_wrap)
338 fprintf(fp, ".nsw");
339 if (instr->no_unsigned_wrap)
340 fprintf(fp, ".nuw");
341 fprintf(fp, " ");
342
343 for (unsigned i = 0; i < nir_op_infos[instr->op].num_inputs; i++) {
344 if (i != 0)
345 fprintf(fp, ", ");
346
347 print_alu_src(instr, i, state);
348 }
349 }
350
351 static const char *
get_var_name(nir_variable * var,print_state * state)352 get_var_name(nir_variable *var, print_state *state)
353 {
354 if (state->ht == NULL)
355 return var->name ? var->name : "unnamed";
356
357 assert(state->syms);
358
359 struct hash_entry *entry = _mesa_hash_table_search(state->ht, var);
360 if (entry)
361 return entry->data;
362
363 char *name;
364 if (var->name == NULL) {
365 name = ralloc_asprintf(state->syms, "@%u", state->index++);
366 } else {
367 struct set_entry *set_entry = _mesa_set_search(state->syms, var->name);
368 if (set_entry != NULL) {
369 /* we have a collision with another name, append an @ + a unique
370 * index */
371 name = ralloc_asprintf(state->syms, "%s@%u", var->name,
372 state->index++);
373 } else {
374 /* Mark this one as seen */
375 _mesa_set_add(state->syms, var->name);
376 name = var->name;
377 }
378 }
379
380 _mesa_hash_table_insert(state->ht, var, name);
381
382 return name;
383 }
384
385 static const char *
get_constant_sampler_addressing_mode(enum cl_sampler_addressing_mode mode)386 get_constant_sampler_addressing_mode(enum cl_sampler_addressing_mode mode)
387 {
388 switch (mode) {
389 case SAMPLER_ADDRESSING_MODE_NONE: return "none";
390 case SAMPLER_ADDRESSING_MODE_CLAMP_TO_EDGE: return "clamp_to_edge";
391 case SAMPLER_ADDRESSING_MODE_CLAMP: return "clamp";
392 case SAMPLER_ADDRESSING_MODE_REPEAT: return "repeat";
393 case SAMPLER_ADDRESSING_MODE_REPEAT_MIRRORED: return "repeat_mirrored";
394 default: unreachable("Invalid addressing mode");
395 }
396 }
397
398 static const char *
get_constant_sampler_filter_mode(enum cl_sampler_filter_mode mode)399 get_constant_sampler_filter_mode(enum cl_sampler_filter_mode mode)
400 {
401 switch (mode) {
402 case SAMPLER_FILTER_MODE_NEAREST: return "nearest";
403 case SAMPLER_FILTER_MODE_LINEAR: return "linear";
404 default: unreachable("Invalid filter mode");
405 }
406 }
407
408 static void
print_constant(nir_constant * c,const struct glsl_type * type,print_state * state)409 print_constant(nir_constant *c, const struct glsl_type *type, print_state *state)
410 {
411 FILE *fp = state->fp;
412 const unsigned rows = glsl_get_vector_elements(type);
413 const unsigned cols = glsl_get_matrix_columns(type);
414 unsigned i;
415
416 switch (glsl_get_base_type(type)) {
417 case GLSL_TYPE_BOOL:
418 /* Only float base types can be matrices. */
419 assert(cols == 1);
420
421 for (i = 0; i < rows; i++) {
422 if (i > 0) fprintf(fp, ", ");
423 fprintf(fp, "%s", c->values[i].b ? "true" : "false");
424 }
425 break;
426
427 case GLSL_TYPE_UINT8:
428 case GLSL_TYPE_INT8:
429 /* Only float base types can be matrices. */
430 assert(cols == 1);
431
432 for (i = 0; i < rows; i++) {
433 if (i > 0) fprintf(fp, ", ");
434 fprintf(fp, "0x%02x", c->values[i].u8);
435 }
436 break;
437
438 case GLSL_TYPE_UINT16:
439 case GLSL_TYPE_INT16:
440 /* Only float base types can be matrices. */
441 assert(cols == 1);
442
443 for (i = 0; i < rows; i++) {
444 if (i > 0) fprintf(fp, ", ");
445 fprintf(fp, "0x%04x", c->values[i].u16);
446 }
447 break;
448
449 case GLSL_TYPE_UINT:
450 case GLSL_TYPE_INT:
451 /* Only float base types can be matrices. */
452 assert(cols == 1);
453
454 for (i = 0; i < rows; i++) {
455 if (i > 0) fprintf(fp, ", ");
456 fprintf(fp, "0x%08x", c->values[i].u32);
457 }
458 break;
459
460 case GLSL_TYPE_FLOAT16:
461 case GLSL_TYPE_FLOAT:
462 case GLSL_TYPE_DOUBLE:
463 if (cols > 1) {
464 for (i = 0; i < cols; i++) {
465 if (i > 0) fprintf(fp, ", ");
466 print_constant(c->elements[i], glsl_get_column_type(type), state);
467 }
468 } else {
469 switch (glsl_get_base_type(type)) {
470 case GLSL_TYPE_FLOAT16:
471 for (i = 0; i < rows; i++) {
472 if (i > 0) fprintf(fp, ", ");
473 fprintf(fp, "%f", _mesa_half_to_float(c->values[i].u16));
474 }
475 break;
476
477 case GLSL_TYPE_FLOAT:
478 for (i = 0; i < rows; i++) {
479 if (i > 0) fprintf(fp, ", ");
480 fprintf(fp, "%f", c->values[i].f32);
481 }
482 break;
483
484 case GLSL_TYPE_DOUBLE:
485 for (i = 0; i < rows; i++) {
486 if (i > 0) fprintf(fp, ", ");
487 fprintf(fp, "%f", c->values[i].f64);
488 }
489 break;
490
491 default:
492 unreachable("Cannot get here from the first level switch");
493 }
494 }
495 break;
496
497 case GLSL_TYPE_UINT64:
498 case GLSL_TYPE_INT64:
499 /* Only float base types can be matrices. */
500 assert(cols == 1);
501
502 for (i = 0; i < cols; i++) {
503 if (i > 0) fprintf(fp, ", ");
504 fprintf(fp, "0x%08" PRIx64, c->values[i].u64);
505 }
506 break;
507
508 case GLSL_TYPE_STRUCT:
509 case GLSL_TYPE_INTERFACE:
510 for (i = 0; i < c->num_elements; i++) {
511 if (i > 0) fprintf(fp, ", ");
512 fprintf(fp, "{ ");
513 print_constant(c->elements[i], glsl_get_struct_field(type, i), state);
514 fprintf(fp, " }");
515 }
516 break;
517
518 case GLSL_TYPE_ARRAY:
519 for (i = 0; i < c->num_elements; i++) {
520 if (i > 0) fprintf(fp, ", ");
521 fprintf(fp, "{ ");
522 print_constant(c->elements[i], glsl_get_array_element(type), state);
523 fprintf(fp, " }");
524 }
525 break;
526
527 default:
528 unreachable("not reached");
529 }
530 }
531
532 static const char *
get_variable_mode_str(nir_variable_mode mode,bool want_local_global_mode)533 get_variable_mode_str(nir_variable_mode mode, bool want_local_global_mode)
534 {
535 switch (mode) {
536 case nir_var_shader_in:
537 return "shader_in";
538 case nir_var_shader_out:
539 return "shader_out";
540 case nir_var_uniform:
541 return "uniform";
542 case nir_var_mem_ubo:
543 return "ubo";
544 case nir_var_system_value:
545 return "system";
546 case nir_var_mem_ssbo:
547 return "ssbo";
548 case nir_var_mem_shared:
549 return "shared";
550 case nir_var_mem_global:
551 return "global";
552 case nir_var_mem_push_const:
553 return "push_const";
554 case nir_var_mem_constant:
555 return "constant";
556 case nir_var_image:
557 return "image";
558 case nir_var_shader_temp:
559 return want_local_global_mode ? "shader_temp" : "";
560 case nir_var_function_temp:
561 return want_local_global_mode ? "function_temp" : "";
562 case nir_var_shader_call_data:
563 return "shader_call_data";
564 case nir_var_ray_hit_attrib:
565 return "ray_hit_attrib";
566 case nir_var_mem_task_payload:
567 return "task_payload";
568 default:
569 if (mode && (mode & nir_var_mem_generic) == mode)
570 return "generic";
571 return "";
572 }
573 }
574
575 static void
print_var_decl(nir_variable * var,print_state * state)576 print_var_decl(nir_variable *var, print_state *state)
577 {
578 FILE *fp = state->fp;
579
580 fprintf(fp, "decl_var ");
581
582 const char *const bindless = (var->data.bindless) ? "bindless " : "";
583 const char *const cent = (var->data.centroid) ? "centroid " : "";
584 const char *const samp = (var->data.sample) ? "sample " : "";
585 const char *const patch = (var->data.patch) ? "patch " : "";
586 const char *const inv = (var->data.invariant) ? "invariant " : "";
587 const char *const per_view = (var->data.per_view) ? "per_view " : "";
588 const char *const per_primitive = (var->data.per_primitive) ? "per_primitive " : "";
589 const char *const ray_query = (var->data.ray_query) ? "ray_query " : "";
590 fprintf(fp, "%s%s%s%s%s%s%s%s%s %s ",
591 bindless, cent, samp, patch, inv, per_view, per_primitive, ray_query,
592 get_variable_mode_str(var->data.mode, false),
593 glsl_interp_mode_name(var->data.interpolation));
594
595 enum gl_access_qualifier access = var->data.access;
596 const char *const coher = (access & ACCESS_COHERENT) ? "coherent " : "";
597 const char *const volat = (access & ACCESS_VOLATILE) ? "volatile " : "";
598 const char *const restr = (access & ACCESS_RESTRICT) ? "restrict " : "";
599 const char *const ronly = (access & ACCESS_NON_WRITEABLE) ? "readonly " : "";
600 const char *const wonly = (access & ACCESS_NON_READABLE) ? "writeonly " : "";
601 const char *const reorder = (access & ACCESS_CAN_REORDER) ? "reorderable " : "";
602 const char *const stream_cache_policy = (access & ACCESS_STREAM_CACHE_POLICY) ?
603 "stream-cache-policy " : "";
604 const char *const include_helpers = (access & ACCESS_INCLUDE_HELPERS) ?
605 "include-helpers " : "";
606 fprintf(fp, "%s%s%s%s%s%s%s%s", coher, volat, restr, ronly, wonly, reorder,
607 stream_cache_policy, include_helpers);
608
609 if (glsl_get_base_type(glsl_without_array(var->type)) == GLSL_TYPE_IMAGE) {
610 fprintf(fp, "%s ", util_format_short_name(var->data.image.format));
611 }
612
613 if (var->data.precision) {
614 const char *precisions[] = {
615 "",
616 "highp",
617 "mediump",
618 "lowp",
619 };
620 fprintf(fp, "%s ", precisions[var->data.precision]);
621 }
622
623 fprintf(fp, "%s %s", glsl_get_type_name(var->type),
624 get_var_name(var, state));
625
626 if (var->data.mode & (nir_var_shader_in |
627 nir_var_shader_out |
628 nir_var_uniform |
629 nir_var_mem_ubo |
630 nir_var_mem_ssbo |
631 nir_var_image)) {
632 const char *loc = NULL;
633 char buf[4];
634
635 switch (state->shader->info.stage) {
636 case MESA_SHADER_VERTEX:
637 if (var->data.mode == nir_var_shader_in)
638 loc = gl_vert_attrib_name(var->data.location);
639 else if (var->data.mode == nir_var_shader_out)
640 loc = gl_varying_slot_name_for_stage(var->data.location,
641 state->shader->info.stage);
642 break;
643 case MESA_SHADER_TASK:
644 case MESA_SHADER_MESH:
645 case MESA_SHADER_GEOMETRY:
646 if ((var->data.mode == nir_var_shader_in) ||
647 (var->data.mode == nir_var_shader_out)) {
648 loc = gl_varying_slot_name_for_stage(var->data.location,
649 state->shader->info.stage);
650 }
651 break;
652 case MESA_SHADER_FRAGMENT:
653 if (var->data.mode == nir_var_shader_in) {
654 loc = gl_varying_slot_name_for_stage(var->data.location,
655 state->shader->info.stage);
656 } else if (var->data.mode == nir_var_shader_out) {
657 loc = gl_frag_result_name(var->data.location);
658 }
659 break;
660 case MESA_SHADER_TESS_CTRL:
661 case MESA_SHADER_TESS_EVAL:
662 case MESA_SHADER_COMPUTE:
663 case MESA_SHADER_KERNEL:
664 default:
665 /* TODO */
666 break;
667 }
668
669 if (!loc) {
670 if (var->data.location == ~0) {
671 loc = "~0";
672 } else {
673 snprintf(buf, sizeof(buf), "%u", var->data.location);
674 loc = buf;
675 }
676 }
677
678 /* For shader I/O vars that have been split to components or packed,
679 * print the fractional location within the input/output.
680 */
681 unsigned int num_components =
682 glsl_get_components(glsl_without_array(var->type));
683 const char *components = NULL;
684 char components_local[18] = {'.' /* the rest is 0-filled */};
685 switch (var->data.mode) {
686 case nir_var_shader_in:
687 case nir_var_shader_out:
688 if (num_components < 16 && num_components != 0) {
689 const char *xyzw = comp_mask_string(num_components);
690 for (int i = 0; i < num_components; i++)
691 components_local[i + 1] = xyzw[i + var->data.location_frac];
692
693 components = components_local;
694 }
695 break;
696 default:
697 break;
698 }
699
700 fprintf(fp, " (%s%s, %u, %u)%s", loc,
701 components ? components : "",
702 var->data.driver_location, var->data.binding,
703 var->data.compact ? " compact" : "");
704 }
705
706 if (var->constant_initializer) {
707 fprintf(fp, " = { ");
708 print_constant(var->constant_initializer, var->type, state);
709 fprintf(fp, " }");
710 }
711 if (glsl_type_is_sampler(var->type) && var->data.sampler.is_inline_sampler) {
712 fprintf(fp, " = { %s, %s, %s }",
713 get_constant_sampler_addressing_mode(var->data.sampler.addressing_mode),
714 var->data.sampler.normalized_coordinates ? "true" : "false",
715 get_constant_sampler_filter_mode(var->data.sampler.filter_mode));
716 }
717 if (var->pointer_initializer)
718 fprintf(fp, " = &%s", get_var_name(var->pointer_initializer, state));
719
720 fprintf(fp, "\n");
721 print_annotation(state, var);
722 }
723
724 static void
print_deref_link(const nir_deref_instr * instr,bool whole_chain,print_state * state)725 print_deref_link(const nir_deref_instr *instr, bool whole_chain, print_state *state)
726 {
727 FILE *fp = state->fp;
728
729 if (instr->deref_type == nir_deref_type_var) {
730 fprintf(fp, "%s", get_var_name(instr->var, state));
731 return;
732 } else if (instr->deref_type == nir_deref_type_cast) {
733 fprintf(fp, "(%s *)", glsl_get_type_name(instr->type));
734 print_src(&instr->parent, state);
735 return;
736 }
737
738 assert(instr->parent.is_ssa);
739 nir_deref_instr *parent =
740 nir_instr_as_deref(instr->parent.ssa->parent_instr);
741
742 /* Is the parent we're going to print a bare cast? */
743 const bool is_parent_cast =
744 whole_chain && parent->deref_type == nir_deref_type_cast;
745
746 /* If we're not printing the whole chain, the parent we print will be a SSA
747 * value that represents a pointer. The only deref type that naturally
748 * gives a pointer is a cast.
749 */
750 const bool is_parent_pointer =
751 !whole_chain || parent->deref_type == nir_deref_type_cast;
752
753 /* Struct derefs have a nice syntax that works on pointers, arrays derefs
754 * do not.
755 */
756 const bool need_deref =
757 is_parent_pointer && instr->deref_type != nir_deref_type_struct;
758
759 /* Cast need extra parens and so * dereferences */
760 if (is_parent_cast || need_deref)
761 fprintf(fp, "(");
762
763 if (need_deref)
764 fprintf(fp, "*");
765
766 if (whole_chain) {
767 print_deref_link(parent, whole_chain, state);
768 } else {
769 print_src(&instr->parent, state);
770 }
771
772 if (is_parent_cast || need_deref)
773 fprintf(fp, ")");
774
775 switch (instr->deref_type) {
776 case nir_deref_type_struct:
777 fprintf(fp, "%s%s", is_parent_pointer ? "->" : ".",
778 glsl_get_struct_elem_name(parent->type, instr->strct.index));
779 break;
780
781 case nir_deref_type_array:
782 case nir_deref_type_ptr_as_array: {
783 if (nir_src_is_const(instr->arr.index)) {
784 fprintf(fp, "[%"PRId64"]", nir_src_as_int(instr->arr.index));
785 } else {
786 fprintf(fp, "[");
787 print_src(&instr->arr.index, state);
788 fprintf(fp, "]");
789 }
790 break;
791 }
792
793 case nir_deref_type_array_wildcard:
794 fprintf(fp, "[*]");
795 break;
796
797 default:
798 unreachable("Invalid deref instruction type");
799 }
800 }
801
802 static void
print_deref_instr(nir_deref_instr * instr,print_state * state)803 print_deref_instr(nir_deref_instr *instr, print_state *state)
804 {
805 FILE *fp = state->fp;
806
807 print_dest(&instr->dest, state);
808
809 switch (instr->deref_type) {
810 case nir_deref_type_var:
811 fprintf(fp, " = deref_var ");
812 break;
813 case nir_deref_type_array:
814 case nir_deref_type_array_wildcard:
815 fprintf(fp, " = deref_array ");
816 break;
817 case nir_deref_type_struct:
818 fprintf(fp, " = deref_struct ");
819 break;
820 case nir_deref_type_cast:
821 fprintf(fp, " = deref_cast ");
822 break;
823 case nir_deref_type_ptr_as_array:
824 fprintf(fp, " = deref_ptr_as_array ");
825 break;
826 default:
827 unreachable("Invalid deref instruction type");
828 }
829
830 /* Only casts naturally return a pointer type */
831 if (instr->deref_type != nir_deref_type_cast)
832 fprintf(fp, "&");
833
834 print_deref_link(instr, false, state);
835
836 fprintf(fp, " (");
837 unsigned modes = instr->modes;
838 while (modes) {
839 int m = u_bit_scan(&modes);
840 fprintf(fp, "%s%s", get_variable_mode_str(1 << m, true),
841 modes ? "|" : "");
842 }
843 fprintf(fp, " %s) ", glsl_get_type_name(instr->type));
844
845 if (instr->deref_type != nir_deref_type_var &&
846 instr->deref_type != nir_deref_type_cast) {
847 /* Print the entire chain as a comment */
848 fprintf(fp, "/* &");
849 print_deref_link(instr, true, state);
850 fprintf(fp, " */");
851 }
852
853 if (instr->deref_type == nir_deref_type_cast) {
854 fprintf(fp, " /* ptr_stride=%u, align_mul=%u, align_offset=%u */",
855 instr->cast.ptr_stride,
856 instr->cast.align_mul, instr->cast.align_offset);
857 }
858 }
859
860 static const char *
vulkan_descriptor_type_name(VkDescriptorType type)861 vulkan_descriptor_type_name(VkDescriptorType type)
862 {
863 switch (type) {
864 case VK_DESCRIPTOR_TYPE_SAMPLER: return "sampler";
865 case VK_DESCRIPTOR_TYPE_COMBINED_IMAGE_SAMPLER: return "texture+sampler";
866 case VK_DESCRIPTOR_TYPE_SAMPLED_IMAGE: return "texture";
867 case VK_DESCRIPTOR_TYPE_STORAGE_IMAGE: return "image";
868 case VK_DESCRIPTOR_TYPE_UNIFORM_TEXEL_BUFFER: return "texture-buffer";
869 case VK_DESCRIPTOR_TYPE_STORAGE_TEXEL_BUFFER: return "image-buffer";
870 case VK_DESCRIPTOR_TYPE_UNIFORM_BUFFER: return "UBO";
871 case VK_DESCRIPTOR_TYPE_STORAGE_BUFFER: return "SSBO";
872 case VK_DESCRIPTOR_TYPE_UNIFORM_BUFFER_DYNAMIC: return "UBO";
873 case VK_DESCRIPTOR_TYPE_STORAGE_BUFFER_DYNAMIC: return "SSBO";
874 case VK_DESCRIPTOR_TYPE_INPUT_ATTACHMENT: return "input-att";
875 case VK_DESCRIPTOR_TYPE_INLINE_UNIFORM_BLOCK: return "inline-UBO";
876 case VK_DESCRIPTOR_TYPE_ACCELERATION_STRUCTURE_KHR: return "accel-struct";
877 default: return "unknown";
878 }
879 }
880
881 static void
print_alu_type(nir_alu_type type,print_state * state)882 print_alu_type(nir_alu_type type, print_state *state)
883 {
884 FILE *fp = state->fp;
885 unsigned size = nir_alu_type_get_type_size(type);
886 const char *name;
887
888 switch (nir_alu_type_get_base_type(type)) {
889 case nir_type_int: name = "int"; break;
890 case nir_type_uint: name = "uint"; break;
891 case nir_type_bool: name = "bool"; break;
892 case nir_type_float: name = "float"; break;
893 default: name = "invalid";
894 }
895 if (size)
896 fprintf(fp, "%s%u", name, size);
897 else
898 fprintf(fp, "%s", name);
899 }
900
901 static void
print_intrinsic_instr(nir_intrinsic_instr * instr,print_state * state)902 print_intrinsic_instr(nir_intrinsic_instr *instr, print_state *state)
903 {
904 const nir_intrinsic_info *info = &nir_intrinsic_infos[instr->intrinsic];
905 unsigned num_srcs = info->num_srcs;
906 FILE *fp = state->fp;
907
908 if (info->has_dest) {
909 print_dest(&instr->dest, state);
910 fprintf(fp, " = ");
911 }
912
913 fprintf(fp, "intrinsic %s (", info->name);
914
915 for (unsigned i = 0; i < num_srcs; i++) {
916 if (i != 0)
917 fprintf(fp, ", ");
918
919 print_src(&instr->src[i], state);
920 }
921
922 fprintf(fp, ") (");
923
924 for (unsigned i = 0; i < info->num_indices; i++) {
925 unsigned idx = info->indices[i];
926 bool print_raw = true;
927 if (i != 0)
928 fprintf(fp, ", ");
929 switch (idx) {
930 case NIR_INTRINSIC_WRITE_MASK: {
931 /* special case wrmask to show it as a writemask.. */
932 unsigned wrmask = nir_intrinsic_write_mask(instr);
933 fprintf(fp, "wrmask=");
934 for (unsigned i = 0; i < instr->num_components; i++)
935 if ((wrmask >> i) & 1)
936 fprintf(fp, "%c", comp_mask_string(instr->num_components)[i]);
937 break;
938 }
939
940 case NIR_INTRINSIC_REDUCTION_OP: {
941 nir_op reduction_op = nir_intrinsic_reduction_op(instr);
942 fprintf(fp, "reduction_op=%s", nir_op_infos[reduction_op].name);
943 break;
944 }
945
946 case NIR_INTRINSIC_IMAGE_DIM: {
947 static const char *dim_name[] = {
948 [GLSL_SAMPLER_DIM_1D] = "1D",
949 [GLSL_SAMPLER_DIM_2D] = "2D",
950 [GLSL_SAMPLER_DIM_3D] = "3D",
951 [GLSL_SAMPLER_DIM_CUBE] = "Cube",
952 [GLSL_SAMPLER_DIM_RECT] = "Rect",
953 [GLSL_SAMPLER_DIM_BUF] = "Buf",
954 [GLSL_SAMPLER_DIM_MS] = "2D-MSAA",
955 [GLSL_SAMPLER_DIM_SUBPASS] = "Subpass",
956 [GLSL_SAMPLER_DIM_SUBPASS_MS] = "Subpass-MSAA",
957 };
958 enum glsl_sampler_dim dim = nir_intrinsic_image_dim(instr);
959 assert(dim < ARRAY_SIZE(dim_name) && dim_name[dim]);
960 fprintf(fp, "image_dim=%s", dim_name[dim]);
961 break;
962 }
963
964 case NIR_INTRINSIC_IMAGE_ARRAY: {
965 bool array = nir_intrinsic_image_array(instr);
966 fprintf(fp, "image_array=%s", array ? "true" : "false");
967 break;
968 }
969
970 case NIR_INTRINSIC_FORMAT: {
971 enum pipe_format format = nir_intrinsic_format(instr);
972 fprintf(fp, "format=%s", util_format_short_name(format));
973 break;
974 }
975
976 case NIR_INTRINSIC_DESC_TYPE: {
977 VkDescriptorType desc_type = nir_intrinsic_desc_type(instr);
978 fprintf(fp, "desc_type=%s", vulkan_descriptor_type_name(desc_type));
979 break;
980 }
981
982 case NIR_INTRINSIC_SRC_TYPE: {
983 fprintf(fp, "src_type=");
984 print_alu_type(nir_intrinsic_src_type(instr), state);
985 break;
986 }
987
988 case NIR_INTRINSIC_DEST_TYPE: {
989 fprintf(fp, "dest_type=");
990 print_alu_type(nir_intrinsic_dest_type(instr), state);
991 break;
992 }
993
994 case NIR_INTRINSIC_SWIZZLE_MASK: {
995 fprintf(fp, "swizzle_mask=");
996 unsigned mask = nir_intrinsic_swizzle_mask(instr);
997 if (instr->intrinsic == nir_intrinsic_quad_swizzle_amd) {
998 for (unsigned i = 0; i < 4; i++)
999 fprintf(fp, "%d", (mask >> (i * 2) & 3));
1000 } else if (instr->intrinsic == nir_intrinsic_masked_swizzle_amd) {
1001 fprintf(fp, "((id & %d) | %d) ^ %d", mask & 0x1F,
1002 (mask >> 5) & 0x1F,
1003 (mask >> 10) & 0x1F);
1004 } else {
1005 fprintf(fp, "%d", mask);
1006 }
1007 break;
1008 }
1009
1010 case NIR_INTRINSIC_MEMORY_SEMANTICS: {
1011 nir_memory_semantics semantics = nir_intrinsic_memory_semantics(instr);
1012 fprintf(fp, "mem_semantics=");
1013 switch (semantics & (NIR_MEMORY_ACQUIRE | NIR_MEMORY_RELEASE)) {
1014 case 0: fprintf(fp, "NONE"); break;
1015 case NIR_MEMORY_ACQUIRE: fprintf(fp, "ACQ"); break;
1016 case NIR_MEMORY_RELEASE: fprintf(fp, "REL"); break;
1017 default: fprintf(fp, "ACQ|REL"); break;
1018 }
1019 if (semantics & (NIR_MEMORY_MAKE_AVAILABLE)) fprintf(fp, "|AVAILABLE");
1020 if (semantics & (NIR_MEMORY_MAKE_VISIBLE)) fprintf(fp, "|VISIBLE");
1021 break;
1022 }
1023
1024 case NIR_INTRINSIC_MEMORY_MODES: {
1025 fprintf(fp, "mem_modes=");
1026 unsigned int modes = nir_intrinsic_memory_modes(instr);
1027 while (modes) {
1028 nir_variable_mode m = u_bit_scan(&modes);
1029 fprintf(fp, "%s%s", get_variable_mode_str(1 << m, true), modes ? "|" : "");
1030 }
1031 break;
1032 }
1033
1034 case NIR_INTRINSIC_EXECUTION_SCOPE:
1035 case NIR_INTRINSIC_MEMORY_SCOPE: {
1036 fprintf(fp, "%s=", nir_intrinsic_index_names[idx]);
1037 nir_scope scope =
1038 idx == NIR_INTRINSIC_MEMORY_SCOPE ? nir_intrinsic_memory_scope(instr)
1039 : nir_intrinsic_execution_scope(instr);
1040 switch (scope) {
1041 case NIR_SCOPE_NONE: fprintf(fp, "NONE"); break;
1042 case NIR_SCOPE_DEVICE: fprintf(fp, "DEVICE"); break;
1043 case NIR_SCOPE_QUEUE_FAMILY: fprintf(fp, "QUEUE_FAMILY"); break;
1044 case NIR_SCOPE_WORKGROUP: fprintf(fp, "WORKGROUP"); break;
1045 case NIR_SCOPE_SHADER_CALL: fprintf(fp, "SHADER_CALL"); break;
1046 case NIR_SCOPE_SUBGROUP: fprintf(fp, "SUBGROUP"); break;
1047 case NIR_SCOPE_INVOCATION: fprintf(fp, "INVOCATION"); break;
1048 }
1049 break;
1050 }
1051
1052 case NIR_INTRINSIC_IO_SEMANTICS: {
1053 struct nir_io_semantics io = nir_intrinsic_io_semantics(instr);
1054 fprintf(fp, "io location=%u slots=%u", io.location, io.num_slots);
1055
1056 if (io.dual_source_blend_index)
1057 fprintf(fp, " dualsrc");
1058
1059 if (io.fb_fetch_output)
1060 fprintf(fp, " fbfetch");
1061
1062 if (io.per_view)
1063 fprintf(fp, " perview");
1064
1065 if (io.medium_precision)
1066 fprintf(fp, " mediump");
1067
1068 if (io.high_16bits)
1069 fprintf(fp, " high_16bits");
1070
1071 if (io.no_varying)
1072 fprintf(fp, " no_varying");
1073
1074 if (io.no_sysval_output)
1075 fprintf(fp, " no_sysval_output");
1076
1077 if (state->shader &&
1078 state->shader->info.stage == MESA_SHADER_GEOMETRY &&
1079 (instr->intrinsic == nir_intrinsic_store_output ||
1080 instr->intrinsic == nir_intrinsic_store_per_primitive_output ||
1081 instr->intrinsic == nir_intrinsic_store_per_vertex_output)) {
1082 unsigned gs_streams = io.gs_streams;
1083 fprintf(fp, " gs_streams(");
1084 for (unsigned i = 0; i < 4; i++) {
1085 fprintf(fp, "%s%c=%u", i ? " " : "", "xyzw"[i],
1086 (gs_streams >> (i * 2)) & 0x3);
1087 }
1088 fprintf(fp, ")");
1089 }
1090
1091 break;
1092 }
1093
1094 case NIR_INTRINSIC_IO_XFB:
1095 case NIR_INTRINSIC_IO_XFB2: {
1096 /* This prints both IO_XFB and IO_XFB2. */
1097 fprintf(fp, "xfb%s(", idx == NIR_INTRINSIC_IO_XFB ? "" : "2");
1098 bool first = true;
1099 for (unsigned i = 0; i < 2; i++) {
1100 unsigned start_comp = (idx == NIR_INTRINSIC_IO_XFB ? 0 : 2) + i;
1101 nir_io_xfb xfb = start_comp < 2 ? nir_intrinsic_io_xfb(instr) :
1102 nir_intrinsic_io_xfb2(instr);
1103
1104 if (!xfb.out[i].num_components)
1105 continue;
1106
1107 if (!first)
1108 fprintf(fp, ", ");
1109 first = false;
1110
1111 if (xfb.out[i].num_components > 1) {
1112 fprintf(fp, "components=%u..%u",
1113 start_comp, start_comp + xfb.out[i].num_components - 1);
1114 } else {
1115 fprintf(fp, "component=%u", start_comp);
1116 }
1117 fprintf(fp, " buffer=%u offset=%u",
1118 xfb.out[i].buffer, (uint32_t)xfb.out[i].offset * 4);
1119 }
1120 fprintf(fp, ")");
1121 break;
1122 }
1123
1124 case NIR_INTRINSIC_ROUNDING_MODE: {
1125 fprintf(fp, "rounding_mode=");
1126 switch (nir_intrinsic_rounding_mode(instr)) {
1127 case nir_rounding_mode_undef: fprintf(fp, "undef"); break;
1128 case nir_rounding_mode_rtne: fprintf(fp, "rtne"); break;
1129 case nir_rounding_mode_ru: fprintf(fp, "ru"); break;
1130 case nir_rounding_mode_rd: fprintf(fp, "rd"); break;
1131 case nir_rounding_mode_rtz: fprintf(fp, "rtz"); break;
1132 default: fprintf(fp, "unkown"); break;
1133 }
1134 break;
1135 }
1136
1137 default: {
1138 unsigned off = info->index_map[idx] - 1;
1139 fprintf(fp, "%s=%d", nir_intrinsic_index_names[idx], instr->const_index[off]);
1140 print_raw = false;
1141 break;
1142 }
1143 }
1144 if (print_raw)
1145 fprintf(fp, " /*%d*/", instr->const_index[i]);
1146 }
1147 fprintf(fp, ")");
1148
1149 if (!state->shader)
1150 return;
1151
1152 nir_variable_mode var_mode;
1153 switch (instr->intrinsic) {
1154 case nir_intrinsic_load_uniform:
1155 var_mode = nir_var_uniform;
1156 break;
1157 case nir_intrinsic_load_input:
1158 case nir_intrinsic_load_interpolated_input:
1159 case nir_intrinsic_load_per_vertex_input:
1160 var_mode = nir_var_shader_in;
1161 break;
1162 case nir_intrinsic_load_output:
1163 case nir_intrinsic_store_output:
1164 case nir_intrinsic_store_per_vertex_output:
1165 var_mode = nir_var_shader_out;
1166 break;
1167 default:
1168 return;
1169 }
1170
1171 nir_foreach_variable_with_modes(var, state->shader, var_mode) {
1172 if ((var->data.driver_location == nir_intrinsic_base(instr)) &&
1173 (instr->intrinsic == nir_intrinsic_load_uniform ||
1174 (nir_intrinsic_component(instr) >= var->data.location_frac &&
1175 nir_intrinsic_component(instr) <
1176 (var->data.location_frac + glsl_get_components(var->type)))) &&
1177 var->name) {
1178 fprintf(fp, "\t/* %s */", var->name);
1179 break;
1180 }
1181 }
1182 }
1183
1184 static void
print_tex_instr(nir_tex_instr * instr,print_state * state)1185 print_tex_instr(nir_tex_instr *instr, print_state *state)
1186 {
1187 FILE *fp = state->fp;
1188
1189 print_dest(&instr->dest, state);
1190
1191 fprintf(fp, " = (");
1192 print_alu_type(instr->dest_type, state);
1193 fprintf(fp, ")");
1194
1195 switch (instr->op) {
1196 case nir_texop_tex:
1197 fprintf(fp, "tex ");
1198 break;
1199 case nir_texop_txb:
1200 fprintf(fp, "txb ");
1201 break;
1202 case nir_texop_txl:
1203 fprintf(fp, "txl ");
1204 break;
1205 case nir_texop_txd:
1206 fprintf(fp, "txd ");
1207 break;
1208 case nir_texop_txf:
1209 fprintf(fp, "txf ");
1210 break;
1211 case nir_texop_txf_ms:
1212 fprintf(fp, "txf_ms ");
1213 break;
1214 case nir_texop_txf_ms_fb:
1215 fprintf(fp, "txf_ms_fb ");
1216 break;
1217 case nir_texop_txf_ms_mcs_intel:
1218 fprintf(fp, "txf_ms_mcs_intel ");
1219 break;
1220 case nir_texop_txs:
1221 fprintf(fp, "txs ");
1222 break;
1223 case nir_texop_lod:
1224 fprintf(fp, "lod ");
1225 break;
1226 case nir_texop_tg4:
1227 fprintf(fp, "tg4 ");
1228 break;
1229 case nir_texop_query_levels:
1230 fprintf(fp, "query_levels ");
1231 break;
1232 case nir_texop_texture_samples:
1233 fprintf(fp, "texture_samples ");
1234 break;
1235 case nir_texop_samples_identical:
1236 fprintf(fp, "samples_identical ");
1237 break;
1238 case nir_texop_tex_prefetch:
1239 fprintf(fp, "tex (pre-dispatchable) ");
1240 break;
1241 case nir_texop_fragment_fetch_amd:
1242 fprintf(fp, "fragment_fetch_amd ");
1243 break;
1244 case nir_texop_fragment_mask_fetch_amd:
1245 fprintf(fp, "fragment_mask_fetch_amd ");
1246 break;
1247 default:
1248 unreachable("Invalid texture operation");
1249 break;
1250 }
1251
1252 bool has_texture_deref = false, has_sampler_deref = false;
1253 for (unsigned i = 0; i < instr->num_srcs; i++) {
1254 if (i > 0) {
1255 fprintf(fp, ", ");
1256 }
1257
1258 print_src(&instr->src[i].src, state);
1259 fprintf(fp, " ");
1260
1261 switch(instr->src[i].src_type) {
1262 case nir_tex_src_backend1:
1263 fprintf(fp, "(backend1)");
1264 break;
1265 case nir_tex_src_backend2:
1266 fprintf(fp, "(backend2)");
1267 break;
1268 case nir_tex_src_coord:
1269 fprintf(fp, "(coord)");
1270 break;
1271 case nir_tex_src_projector:
1272 fprintf(fp, "(projector)");
1273 break;
1274 case nir_tex_src_comparator:
1275 fprintf(fp, "(comparator)");
1276 break;
1277 case nir_tex_src_offset:
1278 fprintf(fp, "(offset)");
1279 break;
1280 case nir_tex_src_bias:
1281 fprintf(fp, "(bias)");
1282 break;
1283 case nir_tex_src_lod:
1284 fprintf(fp, "(lod)");
1285 break;
1286 case nir_tex_src_min_lod:
1287 fprintf(fp, "(min_lod)");
1288 break;
1289 case nir_tex_src_ms_index:
1290 fprintf(fp, "(ms_index)");
1291 break;
1292 case nir_tex_src_ms_mcs_intel:
1293 fprintf(fp, "(ms_mcs_intel)");
1294 break;
1295 case nir_tex_src_ddx:
1296 fprintf(fp, "(ddx)");
1297 break;
1298 case nir_tex_src_ddy:
1299 fprintf(fp, "(ddy)");
1300 break;
1301 case nir_tex_src_texture_deref:
1302 has_texture_deref = true;
1303 fprintf(fp, "(texture_deref)");
1304 break;
1305 case nir_tex_src_sampler_deref:
1306 has_sampler_deref = true;
1307 fprintf(fp, "(sampler_deref)");
1308 break;
1309 case nir_tex_src_texture_offset:
1310 fprintf(fp, "(texture_offset)");
1311 break;
1312 case nir_tex_src_sampler_offset:
1313 fprintf(fp, "(sampler_offset)");
1314 break;
1315 case nir_tex_src_texture_handle:
1316 fprintf(fp, "(texture_handle)");
1317 break;
1318 case nir_tex_src_sampler_handle:
1319 fprintf(fp, "(sampler_handle)");
1320 break;
1321 case nir_tex_src_plane:
1322 fprintf(fp, "(plane)");
1323 break;
1324
1325 default:
1326 unreachable("Invalid texture source type");
1327 break;
1328 }
1329 }
1330
1331 if (instr->op == nir_texop_tg4) {
1332 fprintf(fp, ", %u (gather_component)", instr->component);
1333 }
1334
1335 if (nir_tex_instr_has_explicit_tg4_offsets(instr)) {
1336 fprintf(fp, ", { (%i, %i)", instr->tg4_offsets[0][0], instr->tg4_offsets[0][1]);
1337 for (unsigned i = 1; i < 4; ++i)
1338 fprintf(fp, ", (%i, %i)", instr->tg4_offsets[i][0],
1339 instr->tg4_offsets[i][1]);
1340 fprintf(fp, " } (offsets)");
1341 }
1342
1343 if (instr->op != nir_texop_txf_ms_fb) {
1344 if (!has_texture_deref) {
1345 fprintf(fp, ", %u (texture)", instr->texture_index);
1346 }
1347
1348 if (!has_sampler_deref) {
1349 fprintf(fp, ", %u (sampler)", instr->sampler_index);
1350 }
1351 }
1352
1353 if (instr->texture_non_uniform) {
1354 fprintf(fp, ", texture non-uniform");
1355 }
1356
1357 if (instr->sampler_non_uniform) {
1358 fprintf(fp, ", sampler non-uniform");
1359 }
1360
1361 if (instr->is_sparse) {
1362 fprintf(fp, ", sparse");
1363 }
1364 }
1365
1366 static void
print_call_instr(nir_call_instr * instr,print_state * state)1367 print_call_instr(nir_call_instr *instr, print_state *state)
1368 {
1369 FILE *fp = state->fp;
1370
1371 fprintf(fp, "call %s ", instr->callee->name);
1372
1373 for (unsigned i = 0; i < instr->num_params; i++) {
1374 if (i != 0)
1375 fprintf(fp, ", ");
1376
1377 print_src(&instr->params[i], state);
1378 }
1379 }
1380
1381 static void
print_jump_instr(nir_jump_instr * instr,print_state * state)1382 print_jump_instr(nir_jump_instr *instr, print_state *state)
1383 {
1384 FILE *fp = state->fp;
1385
1386 switch (instr->type) {
1387 case nir_jump_break:
1388 fprintf(fp, "break");
1389 break;
1390
1391 case nir_jump_continue:
1392 fprintf(fp, "continue");
1393 break;
1394
1395 case nir_jump_return:
1396 fprintf(fp, "return");
1397 break;
1398
1399 case nir_jump_halt:
1400 fprintf(fp, "halt");
1401 break;
1402
1403 case nir_jump_goto:
1404 fprintf(fp, "goto block_%u",
1405 instr->target ? instr->target->index : -1);
1406 break;
1407
1408 case nir_jump_goto_if:
1409 fprintf(fp, "goto block_%u if ",
1410 instr->target ? instr->target->index : -1);
1411 print_src(&instr->condition, state);
1412 fprintf(fp, " else block_%u",
1413 instr->else_target ? instr->else_target->index : -1);
1414 break;
1415 }
1416 }
1417
1418 static void
print_ssa_undef_instr(nir_ssa_undef_instr * instr,print_state * state)1419 print_ssa_undef_instr(nir_ssa_undef_instr* instr, print_state *state)
1420 {
1421 FILE *fp = state->fp;
1422 print_ssa_def(&instr->def, state);
1423 fprintf(fp, " = undefined");
1424 }
1425
1426 static void
print_phi_instr(nir_phi_instr * instr,print_state * state)1427 print_phi_instr(nir_phi_instr *instr, print_state *state)
1428 {
1429 FILE *fp = state->fp;
1430 print_dest(&instr->dest, state);
1431 fprintf(fp, " = phi ");
1432 nir_foreach_phi_src(src, instr) {
1433 if (&src->node != exec_list_get_head(&instr->srcs))
1434 fprintf(fp, ", ");
1435
1436 fprintf(fp, "block_%u: ", src->pred->index);
1437 print_src(&src->src, state);
1438 }
1439 }
1440
1441 static void
print_parallel_copy_instr(nir_parallel_copy_instr * instr,print_state * state)1442 print_parallel_copy_instr(nir_parallel_copy_instr *instr, print_state *state)
1443 {
1444 FILE *fp = state->fp;
1445 nir_foreach_parallel_copy_entry(entry, instr) {
1446 if (&entry->node != exec_list_get_head(&instr->entries))
1447 fprintf(fp, "; ");
1448
1449 print_dest(&entry->dest, state);
1450 fprintf(fp, " = ");
1451 print_src(&entry->src, state);
1452 }
1453 }
1454
1455 static void
print_instr(const nir_instr * instr,print_state * state,unsigned tabs)1456 print_instr(const nir_instr *instr, print_state *state, unsigned tabs)
1457 {
1458 FILE *fp = state->fp;
1459 print_tabs(tabs, fp);
1460
1461 switch (instr->type) {
1462 case nir_instr_type_alu:
1463 print_alu_instr(nir_instr_as_alu(instr), state);
1464 break;
1465
1466 case nir_instr_type_deref:
1467 print_deref_instr(nir_instr_as_deref(instr), state);
1468 break;
1469
1470 case nir_instr_type_call:
1471 print_call_instr(nir_instr_as_call(instr), state);
1472 break;
1473
1474 case nir_instr_type_intrinsic:
1475 print_intrinsic_instr(nir_instr_as_intrinsic(instr), state);
1476 break;
1477
1478 case nir_instr_type_tex:
1479 print_tex_instr(nir_instr_as_tex(instr), state);
1480 break;
1481
1482 case nir_instr_type_load_const:
1483 print_load_const_instr(nir_instr_as_load_const(instr), state);
1484 break;
1485
1486 case nir_instr_type_jump:
1487 print_jump_instr(nir_instr_as_jump(instr), state);
1488 break;
1489
1490 case nir_instr_type_ssa_undef:
1491 print_ssa_undef_instr(nir_instr_as_ssa_undef(instr), state);
1492 break;
1493
1494 case nir_instr_type_phi:
1495 print_phi_instr(nir_instr_as_phi(instr), state);
1496 break;
1497
1498 case nir_instr_type_parallel_copy:
1499 print_parallel_copy_instr(nir_instr_as_parallel_copy(instr), state);
1500 break;
1501
1502 default:
1503 unreachable("Invalid instruction type");
1504 break;
1505 }
1506 }
1507
1508 static void print_cf_node(nir_cf_node *node, print_state *state,
1509 unsigned tabs);
1510
1511 static void
print_block(nir_block * block,print_state * state,unsigned tabs)1512 print_block(nir_block *block, print_state *state, unsigned tabs)
1513 {
1514 FILE *fp = state->fp;
1515
1516 print_tabs(tabs, fp);
1517 fprintf(fp, "block block_%u:\n", block->index);
1518
1519 nir_block **preds = nir_block_get_predecessors_sorted(block, NULL);
1520
1521 print_tabs(tabs, fp);
1522 fprintf(fp, "/* preds: ");
1523 for (unsigned i = 0; i < block->predecessors->entries; i++) {
1524 fprintf(fp, "block_%u ", preds[i]->index);
1525 }
1526 fprintf(fp, "*/\n");
1527
1528 ralloc_free(preds);
1529
1530 nir_foreach_instr(instr, block) {
1531 print_instr(instr, state, tabs);
1532 fprintf(fp, "\n");
1533 print_annotation(state, instr);
1534 }
1535
1536 print_tabs(tabs, fp);
1537 fprintf(fp, "/* succs: ");
1538 for (unsigned i = 0; i < 2; i++)
1539 if (block->successors[i]) {
1540 fprintf(fp, "block_%u ", block->successors[i]->index);
1541 }
1542 fprintf(fp, "*/\n");
1543 }
1544
1545 static void
print_if(nir_if * if_stmt,print_state * state,unsigned tabs)1546 print_if(nir_if *if_stmt, print_state *state, unsigned tabs)
1547 {
1548 FILE *fp = state->fp;
1549
1550 print_tabs(tabs, fp);
1551 fprintf(fp, "if ");
1552 print_src(&if_stmt->condition, state);
1553 fprintf(fp, " {\n");
1554 foreach_list_typed(nir_cf_node, node, node, &if_stmt->then_list) {
1555 print_cf_node(node, state, tabs + 1);
1556 }
1557 print_tabs(tabs, fp);
1558 fprintf(fp, "} else {\n");
1559 foreach_list_typed(nir_cf_node, node, node, &if_stmt->else_list) {
1560 print_cf_node(node, state, tabs + 1);
1561 }
1562 print_tabs(tabs, fp);
1563 fprintf(fp, "}\n");
1564 }
1565
1566 static void
print_loop(nir_loop * loop,print_state * state,unsigned tabs)1567 print_loop(nir_loop *loop, print_state *state, unsigned tabs)
1568 {
1569 FILE *fp = state->fp;
1570
1571 print_tabs(tabs, fp);
1572 fprintf(fp, "loop {\n");
1573 foreach_list_typed(nir_cf_node, node, node, &loop->body) {
1574 print_cf_node(node, state, tabs + 1);
1575 }
1576 print_tabs(tabs, fp);
1577 fprintf(fp, "}\n");
1578 }
1579
1580 static void
print_cf_node(nir_cf_node * node,print_state * state,unsigned int tabs)1581 print_cf_node(nir_cf_node *node, print_state *state, unsigned int tabs)
1582 {
1583 switch (node->type) {
1584 case nir_cf_node_block:
1585 print_block(nir_cf_node_as_block(node), state, tabs);
1586 break;
1587
1588 case nir_cf_node_if:
1589 print_if(nir_cf_node_as_if(node), state, tabs);
1590 break;
1591
1592 case nir_cf_node_loop:
1593 print_loop(nir_cf_node_as_loop(node), state, tabs);
1594 break;
1595
1596 default:
1597 unreachable("Invalid CFG node type");
1598 }
1599 }
1600
1601 static void
print_function_impl(nir_function_impl * impl,print_state * state)1602 print_function_impl(nir_function_impl *impl, print_state *state)
1603 {
1604 FILE *fp = state->fp;
1605
1606 fprintf(fp, "\nimpl %s ", impl->function->name);
1607
1608 fprintf(fp, "{\n");
1609
1610 if (impl->preamble) {
1611 fprintf(fp, "\tpreamble %s\n", impl->preamble->name);
1612 }
1613
1614 nir_foreach_function_temp_variable(var, impl) {
1615 fprintf(fp, "\t");
1616 print_var_decl(var, state);
1617 }
1618
1619 foreach_list_typed(nir_register, reg, node, &impl->registers) {
1620 fprintf(fp, "\t");
1621 print_register_decl(reg, state);
1622 }
1623
1624 nir_index_blocks(impl);
1625
1626 foreach_list_typed(nir_cf_node, node, node, &impl->body) {
1627 print_cf_node(node, state, 1);
1628 }
1629
1630 fprintf(fp, "\tblock block_%u:\n}\n\n", impl->end_block->index);
1631 }
1632
1633 static void
print_function(nir_function * function,print_state * state)1634 print_function(nir_function *function, print_state *state)
1635 {
1636 FILE *fp = state->fp;
1637
1638 fprintf(fp, "decl_function %s (%d params)", function->name,
1639 function->num_params);
1640
1641 fprintf(fp, "\n");
1642
1643 if (function->impl != NULL) {
1644 print_function_impl(function->impl, state);
1645 return;
1646 }
1647 }
1648
1649 static void
init_print_state(print_state * state,nir_shader * shader,FILE * fp)1650 init_print_state(print_state *state, nir_shader *shader, FILE *fp)
1651 {
1652 state->fp = fp;
1653 state->shader = shader;
1654 state->ht = _mesa_pointer_hash_table_create(NULL);
1655 state->syms = _mesa_set_create(NULL, _mesa_hash_string,
1656 _mesa_key_string_equal);
1657 state->index = 0;
1658 }
1659
1660 static void
destroy_print_state(print_state * state)1661 destroy_print_state(print_state *state)
1662 {
1663 _mesa_hash_table_destroy(state->ht, NULL);
1664 _mesa_set_destroy(state->syms, NULL);
1665 }
1666
1667 static const char *
primitive_name(unsigned primitive)1668 primitive_name(unsigned primitive)
1669 {
1670 #define PRIM(X) case SHADER_PRIM_ ## X : return #X
1671 switch (primitive) {
1672 PRIM(POINTS);
1673 PRIM(LINES);
1674 PRIM(LINE_LOOP);
1675 PRIM(LINE_STRIP);
1676 PRIM(TRIANGLES);
1677 PRIM(TRIANGLE_STRIP);
1678 PRIM(TRIANGLE_FAN);
1679 PRIM(QUADS);
1680 PRIM(QUAD_STRIP);
1681 PRIM(POLYGON);
1682 default:
1683 return "UNKNOWN";
1684 }
1685 }
1686
1687
1688 void
nir_print_shader_annotated(nir_shader * shader,FILE * fp,struct hash_table * annotations)1689 nir_print_shader_annotated(nir_shader *shader, FILE *fp,
1690 struct hash_table *annotations)
1691 {
1692 print_state state;
1693 init_print_state(&state, shader, fp);
1694
1695 state.annotations = annotations;
1696
1697 fprintf(fp, "shader: %s\n", gl_shader_stage_name(shader->info.stage));
1698
1699 fprintf(fp, "source_sha1: {");
1700 _mesa_sha1_print(fp, shader->info.source_sha1);
1701 fprintf(fp, "}\n");
1702
1703 if (shader->info.name)
1704 fprintf(fp, "name: %s\n", shader->info.name);
1705
1706 if (shader->info.label)
1707 fprintf(fp, "label: %s\n", shader->info.label);
1708
1709 if (gl_shader_stage_uses_workgroup(shader->info.stage)) {
1710 fprintf(fp, "workgroup-size: %u, %u, %u%s\n",
1711 shader->info.workgroup_size[0],
1712 shader->info.workgroup_size[1],
1713 shader->info.workgroup_size[2],
1714 shader->info.workgroup_size_variable ? " (variable)" : "");
1715 fprintf(fp, "shared-size: %u\n", shader->info.shared_size);
1716 }
1717 if (shader->info.stage == MESA_SHADER_MESH ||
1718 shader->info.stage == MESA_SHADER_TASK) {
1719 fprintf(fp, "task_payload-size: %u\n", shader->info.task_payload_size);
1720 }
1721
1722 fprintf(fp, "inputs: %u\n", shader->num_inputs);
1723 fprintf(fp, "outputs: %u\n", shader->num_outputs);
1724 fprintf(fp, "uniforms: %u\n", shader->num_uniforms);
1725 if (shader->info.num_ubos)
1726 fprintf(fp, "ubos: %u\n", shader->info.num_ubos);
1727 fprintf(fp, "shared: %u\n", shader->info.shared_size);
1728 fprintf(fp, "ray queries: %u\n", shader->info.ray_queries);
1729 if (shader->scratch_size)
1730 fprintf(fp, "scratch: %u\n", shader->scratch_size);
1731 if (shader->constant_data_size)
1732 fprintf(fp, "constants: %u\n", shader->constant_data_size);
1733
1734 if (shader->info.stage == MESA_SHADER_GEOMETRY) {
1735 fprintf(fp, "invocations: %u\n", shader->info.gs.invocations);
1736 fprintf(fp, "vertices in: %u\n", shader->info.gs.vertices_in);
1737 fprintf(fp, "vertices out: %u\n", shader->info.gs.vertices_out);
1738 fprintf(fp, "input primitive: %s\n", primitive_name(shader->info.gs.input_primitive));
1739 fprintf(fp, "output primitive: %s\n", primitive_name(shader->info.gs.output_primitive));
1740 fprintf(fp, "active_stream_mask: 0x%x\n", shader->info.gs.active_stream_mask);
1741 fprintf(fp, "uses_end_primitive: %u\n", shader->info.gs.uses_end_primitive);
1742 } else if (shader->info.stage == MESA_SHADER_MESH) {
1743 fprintf(fp, "output primitive: %s\n", primitive_name(shader->info.mesh.primitive_type));
1744 fprintf(fp, "max primitives out: %u\n", shader->info.mesh.max_primitives_out);
1745 fprintf(fp, "max vertices out: %u\n", shader->info.mesh.max_vertices_out);
1746 }
1747
1748 nir_foreach_variable_in_shader(var, shader)
1749 print_var_decl(var, &state);
1750
1751 foreach_list_typed(nir_function, func, node, &shader->functions) {
1752 print_function(func, &state);
1753 }
1754
1755 destroy_print_state(&state);
1756 }
1757
1758 void
nir_print_shader(nir_shader * shader,FILE * fp)1759 nir_print_shader(nir_shader *shader, FILE *fp)
1760 {
1761 nir_print_shader_annotated(shader, fp, NULL);
1762 fflush(fp);
1763 }
1764
1765 char *
nir_shader_as_str_annotated(nir_shader * nir,struct hash_table * annotations,void * mem_ctx)1766 nir_shader_as_str_annotated(nir_shader *nir, struct hash_table *annotations, void *mem_ctx)
1767 {
1768 char *stream_data = NULL;
1769 size_t stream_size = 0;
1770 struct u_memstream mem;
1771 if (u_memstream_open(&mem, &stream_data, &stream_size)) {
1772 FILE *const stream = u_memstream_get(&mem);
1773 nir_print_shader_annotated(nir, stream, annotations);
1774 u_memstream_close(&mem);
1775 }
1776
1777 char *str = ralloc_size(mem_ctx, stream_size + 1);
1778 memcpy(str, stream_data, stream_size);
1779 str[stream_size] = '\0';
1780
1781 free(stream_data);
1782
1783 return str;
1784 }
1785
1786 char *
nir_shader_as_str(nir_shader * nir,void * mem_ctx)1787 nir_shader_as_str(nir_shader *nir, void *mem_ctx)
1788 {
1789 return nir_shader_as_str_annotated(nir, NULL, mem_ctx);
1790 }
1791
1792 void
nir_print_instr(const nir_instr * instr,FILE * fp)1793 nir_print_instr(const nir_instr *instr, FILE *fp)
1794 {
1795 print_state state = {
1796 .fp = fp,
1797 };
1798 if (instr->block) {
1799 nir_function_impl *impl = nir_cf_node_get_function(&instr->block->cf_node);
1800 state.shader = impl->function->shader;
1801 }
1802
1803 print_instr(instr, &state, 0);
1804
1805 }
1806
1807 void
nir_print_deref(const nir_deref_instr * deref,FILE * fp)1808 nir_print_deref(const nir_deref_instr *deref, FILE *fp)
1809 {
1810 print_state state = {
1811 .fp = fp,
1812 };
1813 print_deref_link(deref, true, &state);
1814 }
1815
nir_log_shader_annotated_tagged(enum mesa_log_level level,const char * tag,nir_shader * shader,struct hash_table * annotations)1816 void nir_log_shader_annotated_tagged(enum mesa_log_level level, const char *tag,
1817 nir_shader *shader, struct hash_table *annotations)
1818 {
1819 char *str = nir_shader_as_str_annotated(shader, annotations, NULL);
1820 _mesa_log_multiline(level, tag, str);
1821 ralloc_free(str);
1822 }
1823