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 38static void 39print_tabs(unsigned num_tabs, FILE *fp) 40{ 41 for (unsigned i = 0; i < num_tabs; i++) 42 fprintf(fp, "\t"); 43} 44 45typedef 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 64static void 65print_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 82static void 83print_register(nir_register *reg, print_state *state) 84{ 85 FILE *fp = state->fp; 86 fprintf(fp, "r%u", reg->index); 87} 88 89static const char *sizes[] = { "error", "vec1", "vec2", "vec3", "vec4", 90 "vec5", "error", "error", "vec8", 91 "error", "error", "error", "error", 92 "error", "error", "error", "vec16"}; 93 94static void 95print_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 105static void 106print_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 118static void 119print_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 183static void 184print_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 195static void 196print_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 208static void print_src(const nir_src *src, print_state *state); 209 210static void 211print_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 225static void 226print_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 240static void 241print_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 249static void 250print_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 258static const char * 259comp_mask_string(unsigned num_components) 260{ 261 return (num_components > 4) ? "abcdefghijklmnop" : "xyzw"; 262} 263 264static void 265print_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 307static void 308print_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 325static void 326print_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 351static const char * 352get_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 385static const char * 386get_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 398static const char * 399get_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 408static void 409print_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 532static const char * 533get_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 575static void 576print_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 724static void 725print_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 802static void 803print_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 860static const char * 861vulkan_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 881static void 882print_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 901static void 902print_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 1184static void 1185print_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 1366static void 1367print_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 1381static void 1382print_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 1418static void 1419print_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 1426static void 1427print_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 1441static void 1442print_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 1455static void 1456print_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 1508static void print_cf_node(nir_cf_node *node, print_state *state, 1509 unsigned tabs); 1510 1511static void 1512print_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 1545static void 1546print_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 1566static void 1567print_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 1580static void 1581print_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 1601static void 1602print_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 1633static void 1634print_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 1649static void 1650init_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 1660static void 1661destroy_print_state(print_state *state) 1662{ 1663 _mesa_hash_table_destroy(state->ht, NULL); 1664 _mesa_set_destroy(state->syms, NULL); 1665} 1666 1667static const char * 1668primitive_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 1688void 1689nir_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 1758void 1759nir_print_shader(nir_shader *shader, FILE *fp) 1760{ 1761 nir_print_shader_annotated(shader, fp, NULL); 1762 fflush(fp); 1763} 1764 1765char * 1766nir_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 1786char * 1787nir_shader_as_str(nir_shader *nir, void *mem_ctx) 1788{ 1789 return nir_shader_as_str_annotated(nir, NULL, mem_ctx); 1790} 1791 1792void 1793nir_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 1807void 1808nir_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 1816void 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