Lines Matching +full:04 +full:- +full:mesh
35 #include "util/mesa-blake3.h"
53 /** map from nir_variable -> printable name */
59 /* an index used to make new non-conflicting names */
89 FILE *fp = state->fp; in print_annotation()
91 if (!state->annotations) in print_annotation()
94 struct hash_entry *entry = _mesa_hash_table_search(state->annotations, obj); in print_annotation()
98 const char *note = entry->data; in print_annotation()
99 _mesa_hash_table_remove(state->annotations, entry); in print_annotation()
113 if (state->shader->info.divergence_analysis_run) in divergence_status()
128 FILE *fp = state->fp; in print_def()
130 …const unsigned ssa_padding = state->max_dest_index ? count_digits(state->max_dest_index) - count_d… in print_def()
132 const unsigned padding = (def->bit_size <= 8) + 1 + ssa_padding; in print_def()
135 divergence_status(state, def->divergent), in print_def()
136 def->bit_size, sizes[def->num_components], in print_def()
137 padding, "", state->def_prefix, def->index); in print_def()
143 const unsigned div = state->shader->info.divergence_analysis_run ? 4 : 0; in calculate_padding_for_no_dest()
146 const unsigned ssa_index = count_digits(state->max_dest_index); in calculate_padding_for_no_dest()
154 FILE *fp = state->fp; in print_no_dest_padding()
156 if (state->padding_for_no_dest) in print_no_dest_padding()
157 fprintf(fp, "%*s", state->padding_for_no_dest, ""); in print_no_dest_padding()
165 fprintf(fp, "0x%016" PRIx64, value->u64); in print_hex_padded_const_value()
168 fprintf(fp, "0x%08x", value->u32); in print_hex_padded_const_value()
171 fprintf(fp, "0x%04x", value->u16); in print_hex_padded_const_value()
174 fprintf(fp, "0x%02x", value->u8); in print_hex_padded_const_value()
186 fprintf(fp, "0x%" PRIx64, value->u64); in print_hex_terse_const_value()
189 fprintf(fp, "0x%x", value->u32); in print_hex_terse_const_value()
192 fprintf(fp, "0x%x", value->u16); in print_hex_terse_const_value()
195 fprintf(fp, "0x%x", value->u8); in print_hex_terse_const_value()
207 fprintf(fp, "%f", value->f64); in print_float_const_value()
210 fprintf(fp, "%f", value->f32); in print_float_const_value()
213 fprintf(fp, "%f", _mesa_half_to_float(value->u16)); in print_float_const_value()
225 fprintf(fp, "%+" PRIi64, value->i64); in print_int_const_value()
228 fprintf(fp, "%+d", value->i32); in print_int_const_value()
231 fprintf(fp, "%+d", value->i16); in print_int_const_value()
234 fprintf(fp, "%+d", value->i8); in print_int_const_value()
246 fprintf(fp, "%" PRIu64, value->u64); in print_uint_const_value()
249 fprintf(fp, "%u", value->u32); in print_uint_const_value()
252 fprintf(fp, "%u", value->u16); in print_uint_const_value()
255 fprintf(fp, "%u", value->u8); in print_uint_const_value()
265 FILE *fp = state->fp; in print_const_from_load()
267 const unsigned bit_size = instr->def.bit_size; in print_const_from_load()
268 const unsigned num_components = instr->def.num_components; in print_const_from_load()
278 fprintf(fp, "%s", instr->value[i].b ? "true" : "false"); in print_const_from_load()
288 const nir_const_value *v = &instr->value[i]; in print_const_from_load()
310 F(&instr->value[i], bit_size, fp); \ in print_const_from_load()
324 const nir_const_value *v = &instr->value[i]; in print_const_from_load()
327 needs_signed |= v->i64 < 0; in print_const_from_load()
328 needs_decimal |= v->u64 >= 10; in print_const_from_load()
331 needs_signed |= v->i32 < 0; in print_const_from_load()
332 needs_decimal |= v->u32 >= 10; in print_const_from_load()
335 needs_signed |= v->i16 < 0; in print_const_from_load()
336 needs_decimal |= v->u16 >= 10; in print_const_from_load()
339 needs_signed |= v->i8 < 0; in print_const_from_load()
340 needs_decimal |= v->u8 >= 10; in print_const_from_load()
347 if (state->int_types) { in print_const_from_load()
348 const unsigned index = instr->def.index; in print_const_from_load()
349 const bool inferred_int = BITSET_TEST(state->int_types, index); in print_const_from_load()
350 const bool inferred_float = BITSET_TEST(state->float_types, index); in print_const_from_load()
384 FILE *fp = state->fp; in print_load_const_instr()
386 print_def(&instr->def, state); in print_load_const_instr()
397 FILE *fp = state->fp; in print_src()
398 fprintf(fp, "%s%u", state->def_prefix, src->ssa->index); in print_src()
399 nir_instr *instr = src->ssa->parent_instr; in print_src()
401 if (instr->type == nir_instr_type_load_const && !NIR_DEBUG(PRINT_NO_INLINE_CONSTS)) { in print_src()
407 if (type == nir_type_invalid && state->int_types) { in print_src()
408 const unsigned index = load_const->def.index; in print_src()
409 const bool inferred_int = BITSET_TEST(state->int_types, index); in print_src()
410 const bool inferred_float = BITSET_TEST(state->float_types, index); in print_src()
434 FILE *fp = state->fp; in print_alu_src()
436 const nir_op_info *info = &nir_op_infos[instr->op]; in print_alu_src()
437 print_src(&instr->src[src].src, state, info->input_types[src]); in print_alu_src()
448 if (instr->src[src].swizzle[i] != i) { in print_alu_src()
454 unsigned live_channels = nir_src_num_components(instr->src[src].src); in print_alu_src()
462 fprintf(fp, "%c", comp_mask_string(live_channels)[instr->src[src].swizzle[i]]); in print_alu_src()
470 FILE *fp = state->fp; in print_alu_instr()
472 print_def(&instr->def, state); in print_alu_instr()
474 fprintf(fp, " = %s", nir_op_infos[instr->op].name); in print_alu_instr()
475 if (instr->exact) in print_alu_instr()
477 if (instr->no_signed_wrap) in print_alu_instr()
479 if (instr->no_unsigned_wrap) in print_alu_instr()
483 for (unsigned i = 0; i < nir_op_infos[instr->op].num_inputs; i++) { in print_alu_instr()
494 if (state->ht == NULL) in get_var_name()
495 return var->name ? var->name : "unnamed"; in get_var_name()
497 assert(state->syms); in get_var_name()
499 struct hash_entry *entry = _mesa_hash_table_search(state->ht, var); in get_var_name()
501 return entry->data; in get_var_name()
504 if (var->name == NULL) { in get_var_name()
505 name = ralloc_asprintf(state->syms, "#%u", state->index++); in get_var_name()
507 struct set_entry *set_entry = _mesa_set_search(state->syms, var->name); in get_var_name()
511 name = ralloc_asprintf(state->syms, "%s#%u", var->name, in get_var_name()
512 state->index++); in get_var_name()
515 _mesa_set_add(state->syms, var->name); in get_var_name()
516 name = var->name; in get_var_name()
520 _mesa_hash_table_insert(state->ht, var, name); in get_var_name()
560 FILE *fp = state->fp; in print_constant()
573 fprintf(fp, "%s", c->values[i].b ? "true" : "false"); in print_constant()
585 fprintf(fp, "0x%02x", c->values[i].u8); in print_constant()
597 fprintf(fp, "0x%04x", c->values[i].u16); in print_constant()
609 fprintf(fp, "0x%08x", c->values[i].u32); in print_constant()
620 print_constant(c->elements[i], glsl_get_column_type(type), state); in print_constant()
628 fprintf(fp, "%f", _mesa_half_to_float(c->values[i].u16)); in print_constant()
636 fprintf(fp, "%f", c->values[i].f32); in print_constant()
644 fprintf(fp, "%f", c->values[i].f64); in print_constant()
662 fprintf(fp, "0x%08" PRIx64, c->values[i].u64); in print_constant()
668 for (i = 0; i < c->num_elements; i++) { in print_constant()
672 print_constant(c->elements[i], glsl_get_struct_field(type, i), state); in print_constant()
678 for (i = 0; i < c->num_elements; i++) { in print_constant()
682 print_constant(c->elements[i], glsl_get_array_element(type), state); in print_constant()
789 fputs("none", state->fp); in print_access()
802 { ACCESS_NON_UNIFORM, "non-uniform" }, in print_access()
804 { ACCESS_NON_TEMPORAL, "non-temporal" }, in print_access()
805 { ACCESS_INCLUDE_HELPERS, "include-helpers" }, in print_access()
806 { ACCESS_IS_SWIZZLED_AMD, "is-swizzled-amd" }, in print_access()
807 { ACCESS_USES_FORMAT_AMD, "uses-format-amd" }, in print_access()
808 { ACCESS_FMASK_LOWERED_AMD, "fmask-lowered-amd" }, in print_access()
810 { ACCESS_CP_GE_COHERENT_AMD, "cp-ge-coherent-amd" }, in print_access()
811 { ACCESS_IN_BOUNDS_AGX, "in-bounds-agx" }, in print_access()
812 { ACCESS_KEEP_SCALAR, "keep-scalar" }, in print_access()
813 { ACCESS_SMEM_AMD, "smem-amd" }, in print_access()
819 fprintf(state->fp, "%s%s", first ? "" : separator, modes[i].name); in print_access()
828 FILE *fp = state->fp; in print_var_decl()
832 const char *const bindless = (var->data.bindless) ? "bindless " : ""; in print_var_decl()
833 const char *const cent = (var->data.centroid) ? "centroid " : ""; in print_var_decl()
834 const char *const samp = (var->data.sample) ? "sample " : ""; in print_var_decl()
835 const char *const patch = (var->data.patch) ? "patch " : ""; in print_var_decl()
836 const char *const inv = (var->data.invariant) ? "invariant " : ""; in print_var_decl()
837 const char *const per_view = (var->data.per_view) ? "per_view " : ""; in print_var_decl()
838 const char *const per_primitive = (var->data.per_primitive) ? "per_primitive " : ""; in print_var_decl()
839 const char *const ray_query = (var->data.ray_query) ? "ray_query " : ""; in print_var_decl()
840 const char *const fb_fetch = var->data.fb_fetch_output ? "fb_fetch_output " : ""; in print_var_decl()
844 get_variable_mode_str(var->data.mode, false), in print_var_decl()
845 glsl_interp_mode_name(var->data.interpolation)); in print_var_decl()
847 print_access(var->data.access, state, " "); in print_var_decl()
850 if (glsl_get_base_type(glsl_without_array(var->type)) == GLSL_TYPE_IMAGE) { in print_var_decl()
851 fprintf(fp, "%s ", util_format_short_name(var->data.image.format)); in print_var_decl()
854 if (var->data.precision) { in print_var_decl()
861 fprintf(fp, "%s ", precisions[var->data.precision]); in print_var_decl()
864 fprintf(fp, "%s %s", glsl_get_type_name(var->type), in print_var_decl()
867 if (var->data.mode & (nir_var_shader_in | in print_var_decl()
875 const char *loc = get_location_str(var->data.location, in print_var_decl()
876 state->shader->info.stage, in print_var_decl()
877 var->data.mode, buf); in print_var_decl()
883 glsl_get_components(glsl_without_array(var->type)); in print_var_decl()
885 char components_local[18] = { '.' /* the rest is 0-filled */ }; in print_var_decl()
886 switch (var->data.mode) { in print_var_decl()
892 components_local[i + 1] = xyzw[i + var->data.location_frac]; in print_var_decl()
901 if (var->data.mode & nir_var_system_value) { in print_var_decl()
906 var->data.driver_location, var->data.binding, in print_var_decl()
907 var->data.compact ? " compact" : ""); in print_var_decl()
911 if (var->constant_initializer) { in print_var_decl()
912 if (var->constant_initializer->is_null_constant) { in print_var_decl()
916 print_constant(var->constant_initializer, var->type, state); in print_var_decl()
920 if (glsl_type_is_sampler(var->type) && var->data.sampler.is_inline_sampler) { in print_var_decl()
922 get_constant_sampler_addressing_mode(var->data.sampler.addressing_mode), in print_var_decl()
923 var->data.sampler.normalized_coordinates ? "true" : "false", in print_var_decl()
924 get_constant_sampler_filter_mode(var->data.sampler.filter_mode)); in print_var_decl()
926 if (var->pointer_initializer) in print_var_decl()
927 fprintf(fp, " = &%s", get_var_name(var->pointer_initializer, state)); in print_var_decl()
936 FILE *fp = state->fp; in print_deref_link()
938 if (instr->deref_type == nir_deref_type_var) { in print_deref_link()
939 fprintf(fp, "%s", get_var_name(instr->var, state)); in print_deref_link()
941 } else if (instr->deref_type == nir_deref_type_cast) { in print_deref_link()
942 fprintf(fp, "(%s *)", glsl_get_type_name(instr->type)); in print_deref_link()
943 print_src(&instr->parent, state, nir_type_invalid); in print_deref_link()
948 nir_instr_as_deref(instr->parent.ssa->parent_instr); in print_deref_link()
952 whole_chain && parent->deref_type == nir_deref_type_cast; in print_deref_link()
959 !whole_chain || parent->deref_type == nir_deref_type_cast; in print_deref_link()
965 is_parent_pointer && instr->deref_type != nir_deref_type_struct; in print_deref_link()
977 print_src(&instr->parent, state, nir_type_invalid); in print_deref_link()
983 switch (instr->deref_type) { in print_deref_link()
985 fprintf(fp, "%s%s", is_parent_pointer ? "->" : ".", in print_deref_link()
986 glsl_get_struct_elem_name(parent->type, instr->strct.index)); in print_deref_link()
991 if (nir_src_is_const(instr->arr.index)) { in print_deref_link()
992 fprintf(fp, "[%" PRId64 "]", nir_src_as_int(instr->arr.index)); in print_deref_link()
995 print_src(&instr->arr.index, state, nir_type_invalid); in print_deref_link()
1013 FILE *fp = state->fp; in print_deref_instr()
1015 print_def(&instr->def, state); in print_deref_instr()
1017 switch (instr->deref_type) { in print_deref_instr()
1039 if (instr->deref_type != nir_deref_type_cast) in print_deref_instr()
1045 unsigned modes = instr->modes; in print_deref_instr()
1051 fprintf(fp, " %s)", glsl_get_type_name(instr->type)); in print_deref_instr()
1053 if (instr->deref_type == nir_deref_type_cast) { in print_deref_instr()
1055 instr->cast.ptr_stride, in print_deref_instr()
1056 instr->cast.align_mul, instr->cast.align_offset); in print_deref_instr()
1059 if (instr->deref_type != nir_deref_type_var && in print_deref_instr()
1060 instr->deref_type != nir_deref_type_cast) { in print_deref_instr()
1080 return "texture-buffer"; in vulkan_descriptor_type_name()
1082 return "image-buffer"; in vulkan_descriptor_type_name()
1092 return "input-att"; in vulkan_descriptor_type_name()
1094 return "inline-UBO"; in vulkan_descriptor_type_name()
1096 return "accel-struct"; in vulkan_descriptor_type_name()
1105 FILE *fp = state->fp; in print_alu_type()
1134 const nir_intrinsic_info *info = &nir_intrinsic_infos[instr->intrinsic]; in print_intrinsic_instr()
1135 unsigned num_srcs = info->num_srcs; in print_intrinsic_instr()
1136 FILE *fp = state->fp; in print_intrinsic_instr()
1138 if (info->has_dest) { in print_intrinsic_instr()
1139 print_def(&instr->def, state); in print_intrinsic_instr()
1145 fprintf(fp, "@%s", info->name); in print_intrinsic_instr()
1153 print_src(&instr->src[i], state, nir_intrinsic_instr_src_type(instr, i)); in print_intrinsic_instr()
1159 for (unsigned i = 0; i < info->num_indices; i++) { in print_intrinsic_instr()
1160 unsigned idx = info->indices[i]; in print_intrinsic_instr()
1170 for (unsigned i = 0; i < instr->num_components; i++) in print_intrinsic_instr()
1172 fprintf(fp, "%c", comp_mask_string(instr->num_components)[i]); in print_intrinsic_instr()
1250 [GLSL_SAMPLER_DIM_MS] = "2D-MSAA", in print_intrinsic_instr()
1252 [GLSL_SAMPLER_DIM_SUBPASS_MS] = "Subpass-MSAA", in print_intrinsic_instr()
1293 if (instr->intrinsic == nir_intrinsic_quad_swizzle_amd) { in print_intrinsic_instr()
1296 } else if (instr->intrinsic == nir_intrinsic_masked_swizzle_amd) { in print_intrinsic_instr()
1349 if (strncmp(name, prefix, sizeof(prefix) - 1) == 0) in print_intrinsic_instr()
1350 name += sizeof(prefix) - 1; in print_intrinsic_instr()
1360 switch (instr->intrinsic) { in print_intrinsic_instr()
1386 state->shader->info.stage, mode, in print_intrinsic_instr()
1424 if (state->shader && in print_intrinsic_instr()
1425 state->shader->info.stage == MESA_SHADER_GEOMETRY && in print_intrinsic_instr()
1426 (instr->intrinsic == nir_intrinsic_store_output || in print_intrinsic_instr()
1427 instr->intrinsic == nir_intrinsic_store_per_primitive_output || in print_intrinsic_instr()
1428 instr->intrinsic == nir_intrinsic_store_per_vertex_output || in print_intrinsic_instr()
1429 instr->intrinsic == nir_intrinsic_store_per_view_output)) { in print_intrinsic_instr()
1460 start_comp, start_comp + xfb.out[i].num_components - 1); in print_intrinsic_instr()
1547 fprintf(fp, "non-uniform"); in print_intrinsic_instr()
1550 fprintf(fp, "sampler-embedded"); in print_intrinsic_instr()
1631 unsigned off = info->index_map[idx] - 1; in print_intrinsic_instr()
1632 fprintf(fp, "%s=%d", nir_intrinsic_index_names[idx], instr->const_index[off]); in print_intrinsic_instr()
1637 if (info->num_indices) in print_intrinsic_instr()
1640 if (!state->shader) in print_intrinsic_instr()
1644 switch (instr->intrinsic) { in print_intrinsic_instr()
1664 if (instr->name) { in print_intrinsic_instr()
1665 fprintf(fp, " // %s", instr->name); in print_intrinsic_instr()
1669 nir_foreach_variable_with_modes(var, state->shader, var_mode) { in print_intrinsic_instr()
1670 if (!var->name) in print_intrinsic_instr()
1674 if (instr->intrinsic == nir_intrinsic_load_uniform) { in print_intrinsic_instr()
1675 match = var->data.driver_location == nir_intrinsic_base(instr); in print_intrinsic_instr()
1677 match = nir_intrinsic_component(instr) >= var->data.location_frac && in print_intrinsic_instr()
1679 (var->data.location_frac + glsl_get_components(var->type)); in print_intrinsic_instr()
1683 fprintf(fp, " // %s", var->name); in print_intrinsic_instr()
1692 FILE *fp = state->fp; in print_tex_instr()
1694 print_def(&instr->def, state); in print_tex_instr()
1697 print_alu_type(instr->dest_type, state); in print_tex_instr()
1700 switch (instr->op) { in print_tex_instr()
1744 fprintf(fp, "tex (pre-dispatchable) "); in print_tex_instr()
1779 for (unsigned i = 0; i < instr->num_srcs; i++) { in print_tex_instr()
1784 print_src(&instr->src[i].src, state, nir_tex_instr_src_type(instr, i)); in print_tex_instr()
1787 switch (instr->src[i].src_type) { in print_tex_instr()
1865 if (instr->is_gather_implicit_lod) in print_tex_instr()
1868 if (instr->op == nir_texop_tg4) { in print_tex_instr()
1869 fprintf(fp, ", %u (gather_component)", instr->component); in print_tex_instr()
1873 fprintf(fp, ", { (%i, %i)", instr->tg4_offsets[0][0], instr->tg4_offsets[0][1]); in print_tex_instr()
1875 fprintf(fp, ", (%i, %i)", instr->tg4_offsets[i][0], in print_tex_instr()
1876 instr->tg4_offsets[i][1]); in print_tex_instr()
1880 if (instr->op != nir_texop_txf_ms_fb && !has_texture_deref) { in print_tex_instr()
1881 fprintf(fp, ", %u (texture)", instr->texture_index); in print_tex_instr()
1885 fprintf(fp, ", %u (sampler)", instr->sampler_index); in print_tex_instr()
1888 if (instr->texture_non_uniform) { in print_tex_instr()
1889 fprintf(fp, ", texture non-uniform"); in print_tex_instr()
1892 if (instr->sampler_non_uniform) { in print_tex_instr()
1893 fprintf(fp, ", sampler non-uniform"); in print_tex_instr()
1896 if (instr->is_sparse) { in print_tex_instr()
1904 FILE *fp = state->fp; in print_call_instr()
1908 bool indirect = instr->indirect_callee.ssa; in print_call_instr()
1910 fprintf(fp, "call %s ", instr->callee->name); in print_call_instr()
1913 print_src(&instr->indirect_callee, state, nir_type_invalid); in print_call_instr()
1917 for (unsigned i = 0; i < instr->num_params; i++) { in print_call_instr()
1921 if (instr->callee->params[i].name) in print_call_instr()
1922 fprintf(fp, "%s ", instr->callee->params[i].name); in print_call_instr()
1924 print_src(&instr->params[i], state, nir_type_invalid); in print_call_instr()
1931 FILE *fp = state->fp; in print_jump_instr()
1935 switch (instr->type) { in print_jump_instr()
1954 instr->target ? instr->target->index : -1); in print_jump_instr()
1959 instr->target ? instr->target->index : -1); in print_jump_instr()
1960 print_src(&instr->condition, state, nir_type_invalid); in print_jump_instr()
1962 instr->else_target ? instr->else_target->index : -1); in print_jump_instr()
1970 FILE *fp = state->fp; in print_ssa_undef_instr()
1971 print_def(&instr->def, state); in print_ssa_undef_instr()
1978 FILE *fp = state->fp; in print_phi_instr()
1979 print_def(&instr->def, state); in print_phi_instr()
1982 if (&src->node != exec_list_get_head(&instr->srcs)) in print_phi_instr()
1985 fprintf(fp, "b%u: ", src->pred->index); in print_phi_instr()
1986 print_src(&src->src, state, nir_type_invalid); in print_phi_instr()
1993 FILE *fp = state->fp; in print_parallel_copy_instr()
1995 if (&entry->node != exec_list_get_head(&instr->entries)) in print_parallel_copy_instr()
1998 if (entry->dest_is_reg) { in print_parallel_copy_instr()
2000 print_src(&entry->dest.reg, state, nir_type_invalid); in print_parallel_copy_instr()
2002 print_def(&entry->dest.def, state); in print_parallel_copy_instr()
2006 if (entry->src_is_reg) in print_parallel_copy_instr()
2008 print_src(&entry->src, state, nir_type_invalid); in print_parallel_copy_instr()
2015 FILE *fp = state->fp; in print_debug_info_instr()
2017 switch (instr->type) { in print_debug_info_instr()
2019 fprintf(fp, "// 0x%x", instr->src_loc.spirv_offset); in print_debug_info_instr()
2020 if (instr->src_loc.line) in print_debug_info_instr()
2021 …fprintf(fp, " %s:%u:%u", nir_src_as_string(instr->src_loc.filename), instr->src_loc.line, instr->s… in print_debug_info_instr()
2033 FILE *fp = state->fp; in print_instr()
2035 if (state->debug_info) { in print_instr()
2036 nir_debug_info_instr *di = state->debug_info[instr->index]; in print_instr()
2038 di->src_loc.column = (uint32_t)ftell(fp); in print_instr()
2043 switch (instr->type) { in print_instr()
2093 if (NIR_DEBUG(PRINT_PASS_FLAGS) && instr->pass_flags) in print_instr()
2094 fprintf(fp, " (pass_flags: 0x%x)", instr->pass_flags); in print_instr()
2101 switch (instr->type) { in block_has_instruction_with_dest()
2113 const nir_intrinsic_info *info = &nir_intrinsic_infos[intrin->intrinsic]; in block_has_instruction_with_dest()
2114 if (info->has_dest) in block_has_instruction_with_dest()
2138 FILE *fp = state->fp; in print_block_preds()
2140 for (unsigned i = 0; i < block->predecessors->entries; i++) { in print_block_preds()
2141 fprintf(fp, " b%u", preds[i]->index); in print_block_preds()
2149 FILE *fp = state->fp; in print_block_succs()
2151 if (block->successors[i]) { in print_block_succs()
2152 fprintf(fp, " b%u", block->successors[i]->index); in print_block_succs()
2160 FILE *fp = state->fp; in print_block()
2163 state->padding_for_no_dest = calculate_padding_for_no_dest(state); in print_block()
2165 state->padding_for_no_dest = 0; in print_block()
2169 divergence_status(state, block->divergent), in print_block()
2170 block->index); in print_block()
2172 const bool empty_block = exec_list_is_empty(&block->instr_list); in print_block()
2182 const unsigned block_length = 7 + count_digits(block->index) + 1; in print_block()
2183 …const unsigned pred_padding = block_length < state->padding_for_no_dest ? state->padding_for_no_de… in print_block()
2196 fprintf(fp, "%*s// succs:", state->padding_for_no_dest, ""); in print_block()
2204 FILE *fp = state->fp; in print_if()
2208 print_src(&if_stmt->condition, state, nir_type_invalid); in print_if()
2209 switch (if_stmt->control) { in print_if()
2224 foreach_list_typed(nir_cf_node, node, node, &if_stmt->then_list) { in print_if()
2229 foreach_list_typed(nir_cf_node, node, node, &if_stmt->else_list) { in print_if()
2239 FILE *fp = state->fp; in print_loop()
2242 fprintf(fp, "%sloop {\n", divergence_status(state, loop->divergent_break)); in print_loop()
2243 foreach_list_typed(nir_cf_node, node, node, &loop->body) { in print_loop()
2250 foreach_list_typed(nir_cf_node, node, node, &loop->continue_list) { in print_loop()
2262 switch (node->type) { in print_cf_node()
2283 FILE *fp = state->fp; in print_function_impl()
2285 state->max_dest_index = impl->ssa_alloc; in print_function_impl()
2287 fprintf(fp, "\nimpl %s ", impl->function->name); in print_function_impl()
2291 if (impl->preamble) { in print_function_impl()
2293 fprintf(fp, "preamble %s\n", impl->preamble->name); in print_function_impl()
2301 state->float_types = calloc(BITSET_WORDS(impl->ssa_alloc), sizeof(BITSET_WORD)); in print_function_impl()
2302 state->int_types = calloc(BITSET_WORDS(impl->ssa_alloc), sizeof(BITSET_WORD)); in print_function_impl()
2303 nir_gather_types(impl, state->float_types, state->int_types); in print_function_impl()
2313 foreach_list_typed(nir_cf_node, node, node, &impl->body) { in print_function_impl()
2318 fprintf(fp, "block b%u:\n}\n\n", impl->end_block->index); in print_function_impl()
2320 free(state->float_types); in print_function_impl()
2321 free(state->int_types); in print_function_impl()
2322 state->max_dest_index = 0; in print_function_impl()
2328 FILE *fp = state->fp; in print_function()
2330 fprintf(fp, "decl_function %s (", function->name); in print_function()
2332 for (unsigned i = 0; i < function->num_params; ++i) { in print_function()
2337 nir_parameter param = function->params[i]; in print_function()
2353 /* clang-format off */ in print_function()
2354 fprintf(fp, "%s%s%s", function->dont_inline ? " (noinline)" : in print_function()
2355 function->should_inline ? " (inline)" : "", in print_function()
2356 function->is_exported ? " (exported)" : "", in print_function()
2357 function->is_entrypoint ? " (entrypoint)" : ""); in print_function()
2358 /* clang-format on */ in print_function()
2360 if (function->workgroup_size[0]) { in print_function()
2362 function->workgroup_size[0], in print_function()
2363 function->workgroup_size[1], in print_function()
2364 function->workgroup_size[2]); in print_function()
2369 if (function->impl != NULL) { in print_function()
2370 print_function_impl(function->impl, state); in print_function()
2378 state->fp = fp; in init_print_state()
2379 state->shader = shader; in init_print_state()
2380 state->ht = _mesa_pointer_hash_table_create(NULL); in init_print_state()
2381 state->syms = _mesa_set_create(NULL, _mesa_hash_string, in init_print_state()
2383 state->index = 0; in init_print_state()
2384 state->int_types = NULL; in init_print_state()
2385 state->float_types = NULL; in init_print_state()
2386 state->max_dest_index = 0; in init_print_state()
2387 state->padding_for_no_dest = 0; in init_print_state()
2393 _mesa_hash_table_destroy(state->ht, NULL); in destroy_print_state()
2394 _mesa_set_destroy(state->syms, NULL); in destroy_print_state()
2425 /* Iterate back-to-front to get proper digit order (most significant first). */ in print_bitset()
2426 for (int i = size - 1; i >= 0; --i) { in print_bitset()
2427 fprintf(fp, (i == size - 1) ? "0x%08x" : "'%08x", words[i]); in print_bitset()
2448 /* Print uint64_t value, only if non-zero.
2450 * E.g. inputs_read: 0,15-17
2465 snprintf(buf, sizeof(buf), is_first ? "%d-%d" : ",%d-%d", start, start + count - 1); in print_nz_x64()
2476 /* Print uint32_t value in hex, only if non-zero */
2484 /* Print uint16_t value in hex, only if non-zero */
2489 fprintf(fp, "%s: 0x%04x\n", label, value); in print_nz_x16()
2492 /* Print uint8_t value in hex, only if non-zero */
2500 /* Print unsigned value in decimal, only if non-zero */
2519 fprintf(fp, "shader: %s\n", gl_shader_stage_name(info->stage)); in print_shader_info()
2521 if (memcmp(info->source_blake3, &(blake3_hash){0}, sizeof(info->source_blake3))) { in print_shader_info()
2523 _mesa_blake3_print(fp, info->source_blake3); in print_shader_info()
2527 if (info->name) in print_shader_info()
2528 fprintf(fp, "name: %s\n", info->name); in print_shader_info()
2530 if (info->label) in print_shader_info()
2531 fprintf(fp, "label: %s\n", info->label); in print_shader_info()
2533 print_nz_bool(fp, "internal", info->internal); in print_shader_info()
2535 if (gl_shader_stage_uses_workgroup(info->stage)) { in print_shader_info()
2537 info->workgroup_size[0], in print_shader_info()
2538 info->workgroup_size[1], in print_shader_info()
2539 info->workgroup_size[2], in print_shader_info()
2540 info->workgroup_size_variable ? " (variable)" : ""); in print_shader_info()
2543 if (info->next_stage != MESA_SHADER_NONE) in print_shader_info()
2544 fprintf(fp, "next_stage: %s\n", gl_shader_stage_name(info->next_stage)); in print_shader_info()
2546 print_nz_unsigned(fp, "num_textures", info->num_textures); in print_shader_info()
2547 print_nz_unsigned(fp, "num_ubos", info->num_ubos); in print_shader_info()
2548 print_nz_unsigned(fp, "num_abos", info->num_abos); in print_shader_info()
2549 print_nz_unsigned(fp, "num_ssbos", info->num_ssbos); in print_shader_info()
2550 print_nz_unsigned(fp, "num_images", info->num_images); in print_shader_info()
2552 print_nz_x64(fp, "inputs_read", info->inputs_read); in print_shader_info()
2553 print_nz_x64(fp, "dual_slot_inputs", info->dual_slot_inputs); in print_shader_info()
2554 print_nz_x64(fp, "outputs_written", info->outputs_written); in print_shader_info()
2555 print_nz_x64(fp, "outputs_read", info->outputs_read); in print_shader_info()
2557 …print_nz_bitset(fp, "system_values_read", info->system_values_read, ARRAY_SIZE(info->system_values… in print_shader_info()
2559 print_nz_x64(fp, "per_primitive_inputs", info->per_primitive_inputs); in print_shader_info()
2560 print_nz_x64(fp, "per_primitive_outputs", info->per_primitive_outputs); in print_shader_info()
2561 print_nz_x64(fp, "per_view_outputs", info->per_view_outputs); in print_shader_info()
2563 print_nz_x16(fp, "inputs_read_16bit", info->inputs_read_16bit); in print_shader_info()
2564 print_nz_x16(fp, "outputs_written_16bit", info->outputs_written_16bit); in print_shader_info()
2565 print_nz_x16(fp, "outputs_read_16bit", info->outputs_read_16bit); in print_shader_info()
2566 print_nz_x16(fp, "inputs_read_indirectly_16bit", info->inputs_read_indirectly_16bit); in print_shader_info()
2567 print_nz_x16(fp, "outputs_accessed_indirectly_16bit", info->outputs_accessed_indirectly_16bit); in print_shader_info()
2569 print_nz_x32(fp, "patch_inputs_read", info->patch_inputs_read); in print_shader_info()
2570 print_nz_x32(fp, "patch_outputs_written", info->patch_outputs_written); in print_shader_info()
2571 print_nz_x32(fp, "patch_outputs_read", info->patch_outputs_read); in print_shader_info()
2573 print_nz_x64(fp, "inputs_read_indirectly", info->inputs_read_indirectly); in print_shader_info()
2574 print_nz_x64(fp, "outputs_accessed_indirectly", info->outputs_accessed_indirectly); in print_shader_info()
2575 print_nz_x64(fp, "patch_inputs_read_indirectly", info->patch_inputs_read_indirectly); in print_shader_info()
2576 print_nz_x64(fp, "patch_outputs_accessed_indirectly", info->patch_outputs_accessed_indirectly); in print_shader_info()
2578 print_nz_bitset(fp, "textures_used", info->textures_used, ARRAY_SIZE(info->textures_used)); in print_shader_info()
2579 …print_nz_bitset(fp, "textures_used_by_txf", info->textures_used_by_txf, ARRAY_SIZE(info->textures_… in print_shader_info()
2580 print_nz_bitset(fp, "samplers_used", info->samplers_used, ARRAY_SIZE(info->samplers_used)); in print_shader_info()
2581 print_nz_bitset(fp, "images_used", info->images_used, ARRAY_SIZE(info->images_used)); in print_shader_info()
2582 print_nz_bitset(fp, "image_buffers", info->image_buffers, ARRAY_SIZE(info->image_buffers)); in print_shader_info()
2583 print_nz_bitset(fp, "msaa_images", info->msaa_images, ARRAY_SIZE(info->msaa_images)); in print_shader_info()
2585 print_nz_x32(fp, "float_controls_execution_mode", info->float_controls_execution_mode); in print_shader_info()
2587 print_nz_unsigned(fp, "shared_size", info->shared_size); in print_shader_info()
2589 if (info->stage == MESA_SHADER_MESH || info->stage == MESA_SHADER_TASK) { in print_shader_info()
2590 fprintf(fp, "task_payload_size: %u\n", info->task_payload_size); in print_shader_info()
2593 print_nz_unsigned(fp, "ray queries", info->ray_queries); in print_shader_info()
2595 fprintf(fp, "subgroup_size: %u\n", info->subgroup_size); in print_shader_info()
2597 print_nz_bool(fp, "uses_wide_subgroup_intrinsics", info->uses_wide_subgroup_intrinsics); in print_shader_info()
2599 …bool has_xfb_stride = info->xfb_stride[0] || info->xfb_stride[1] || info->xfb_stride[2] || info->x… in print_shader_info()
2602 info->xfb_stride[0], in print_shader_info()
2603 info->xfb_stride[1], in print_shader_info()
2604 info->xfb_stride[2], in print_shader_info()
2605 info->xfb_stride[3]); in print_shader_info()
2607 …offsets = info->inlinable_uniform_dw_offsets[0] || info->inlinable_uniform_dw_offsets[1] || info->… in print_shader_info()
2610 info->inlinable_uniform_dw_offsets[0], in print_shader_info()
2611 info->inlinable_uniform_dw_offsets[1], in print_shader_info()
2612 info->inlinable_uniform_dw_offsets[2], in print_shader_info()
2613 info->inlinable_uniform_dw_offsets[3]); in print_shader_info()
2615 print_nz_unsigned(fp, "num_inlinable_uniforms", info->num_inlinable_uniforms); in print_shader_info()
2616 print_nz_unsigned(fp, "clip_distance_array_size", info->clip_distance_array_size); in print_shader_info()
2617 print_nz_unsigned(fp, "cull_distance_array_size", info->cull_distance_array_size); in print_shader_info()
2619 print_nz_bool(fp, "uses_texture_gather", info->uses_texture_gather); in print_shader_info()
2620 print_nz_bool(fp, "uses_resource_info_query", info->uses_resource_info_query); in print_shader_info()
2621 print_nz_bool(fp, "divergence_analysis_run", info->divergence_analysis_run); in print_shader_info()
2623 print_nz_x8(fp, "bit_sizes_float", info->bit_sizes_float); in print_shader_info()
2624 print_nz_x8(fp, "bit_sizes_int", info->bit_sizes_int); in print_shader_info()
2626 print_nz_bool(fp, "first_ubo_is_default_ubo", info->first_ubo_is_default_ubo); in print_shader_info()
2627 print_nz_bool(fp, "separate_shader", info->separate_shader); in print_shader_info()
2628 print_nz_bool(fp, "has_transform_feedback_varyings", info->has_transform_feedback_varyings); in print_shader_info()
2629 print_nz_bool(fp, "flrp_lowered", info->flrp_lowered); in print_shader_info()
2630 print_nz_bool(fp, "io_lowered", info->io_lowered); in print_shader_info()
2631 print_nz_bool(fp, "writes_memory", info->writes_memory); in print_shader_info()
2632 print_nz_unsigned(fp, "derivative_group", info->derivative_group); in print_shader_info()
2634 switch (info->stage) { in print_shader_info()
2636 print_nz_x64(fp, "double_inputs", info->vs.double_inputs); in print_shader_info()
2637 print_nz_unsigned(fp, "blit_sgprs_amd", info->vs.blit_sgprs_amd); in print_shader_info()
2638 print_nz_bool(fp, "window_space_position", info->vs.window_space_position); in print_shader_info()
2639 print_nz_bool(fp, "needs_edge_flag", info->vs.needs_edge_flag); in print_shader_info()
2644 fprintf(fp, "primitive_mode: %u\n", info->tess._primitive_mode); in print_shader_info()
2645 fprintf(fp, "tcs_vertices_out: %u\n", info->tess.tcs_vertices_out); in print_shader_info()
2646 fprintf(fp, "spacing: %u\n", info->tess.spacing); in print_shader_info()
2648 print_nz_bool(fp, "ccw", info->tess.ccw); in print_shader_info()
2649 print_nz_bool(fp, "point_mode", info->tess.point_mode); in print_shader_info()
2651 info->tess.tcs_same_invocation_inputs_read); in print_shader_info()
2652 … print_nz_x64(fp, "tcs_cross_invocation_inputs_read", info->tess.tcs_cross_invocation_inputs_read); in print_shader_info()
2653 …print_nz_x64(fp, "tcs_cross_invocation_outputs_read", info->tess.tcs_cross_invocation_outputs_read… in print_shader_info()
2657 fprintf(fp, "output_primitive: %s\n", primitive_name(info->gs.output_primitive)); in print_shader_info()
2658 fprintf(fp, "input_primitive: %s\n", primitive_name(info->gs.input_primitive)); in print_shader_info()
2659 fprintf(fp, "vertices_out: %u\n", info->gs.vertices_out); in print_shader_info()
2660 fprintf(fp, "invocations: %u\n", info->gs.invocations); in print_shader_info()
2661 fprintf(fp, "vertices_in: %u\n", info->gs.vertices_in); in print_shader_info()
2662 print_nz_bool(fp, "uses_end_primitive", info->gs.uses_end_primitive); in print_shader_info()
2663 fprintf(fp, "active_stream_mask: 0x%02x\n", info->gs.active_stream_mask); in print_shader_info()
2667 print_nz_bool(fp, "uses_discard", info->fs.uses_discard); in print_shader_info()
2668 print_nz_bool(fp, "uses_fbfetch_output", info->fs.uses_fbfetch_output); in print_shader_info()
2669 print_nz_bool(fp, "color_is_dual_source", info->fs.color_is_dual_source); in print_shader_info()
2671 print_nz_bool(fp, "require_full_quads", info->fs.require_full_quads); in print_shader_info()
2672 print_nz_bool(fp, "needs_quad_helper_invocations", info->fs.needs_quad_helper_invocations); in print_shader_info()
2673 print_nz_bool(fp, "uses_sample_qualifier", info->fs.uses_sample_qualifier); in print_shader_info()
2674 print_nz_bool(fp, "uses_sample_shading", info->fs.uses_sample_shading); in print_shader_info()
2675 print_nz_bool(fp, "early_fragment_tests", info->fs.early_fragment_tests); in print_shader_info()
2676 print_nz_bool(fp, "inner_coverage", info->fs.inner_coverage); in print_shader_info()
2677 print_nz_bool(fp, "post_depth_coverage", info->fs.post_depth_coverage); in print_shader_info()
2679 print_nz_bool(fp, "pixel_center_integer", info->fs.pixel_center_integer); in print_shader_info()
2680 print_nz_bool(fp, "origin_upper_left", info->fs.origin_upper_left); in print_shader_info()
2681 print_nz_bool(fp, "pixel_interlock_ordered", info->fs.pixel_interlock_ordered); in print_shader_info()
2682 print_nz_bool(fp, "pixel_interlock_unordered", info->fs.pixel_interlock_unordered); in print_shader_info()
2683 print_nz_bool(fp, "sample_interlock_ordered", info->fs.sample_interlock_ordered); in print_shader_info()
2684 print_nz_bool(fp, "sample_interlock_unordered", info->fs.sample_interlock_unordered); in print_shader_info()
2685 print_nz_bool(fp, "untyped_color_outputs", info->fs.untyped_color_outputs); in print_shader_info()
2687 print_nz_unsigned(fp, "depth_layout", info->fs.depth_layout); in print_shader_info()
2689 if (info->fs.color0_interp != INTERP_MODE_NONE) { in print_shader_info()
2691 glsl_interp_mode_name(info->fs.color0_interp)); in print_shader_info()
2693 print_nz_bool(fp, "color0_sample", info->fs.color0_sample); in print_shader_info()
2694 print_nz_bool(fp, "color0_centroid", info->fs.color0_centroid); in print_shader_info()
2696 if (info->fs.color1_interp != INTERP_MODE_NONE) { in print_shader_info()
2698 glsl_interp_mode_name(info->fs.color1_interp)); in print_shader_info()
2700 print_nz_bool(fp, "color1_sample", info->fs.color1_sample); in print_shader_info()
2701 print_nz_bool(fp, "color1_centroid", info->fs.color1_centroid); in print_shader_info()
2703 print_nz_x32(fp, "advanced_blend_modes", info->fs.advanced_blend_modes); in print_shader_info()
2708 …if (info->cs.workgroup_size_hint[0] || info->cs.workgroup_size_hint[1] || info->cs.workgroup_size_… in print_shader_info()
2710 info->cs.workgroup_size_hint[0], in print_shader_info()
2711 info->cs.workgroup_size_hint[1], in print_shader_info()
2712 info->cs.workgroup_size_hint[2]); in print_shader_info()
2713 print_nz_unsigned(fp, "user_data_components_amd", info->cs.user_data_components_amd); in print_shader_info()
2714 print_nz_unsigned(fp, "ptr_size", info->cs.ptr_size); in print_shader_info()
2718 …print_nz_x64(fp, "ms_cross_invocation_output_access", info->mesh.ms_cross_invocation_output_access… in print_shader_info()
2719 fprintf(fp, "max_vertices_out: %u\n", info->mesh.max_vertices_out); in print_shader_info()
2720 fprintf(fp, "max_primitives_out: %u\n", info->mesh.max_primitives_out); in print_shader_info()
2721 fprintf(fp, "primitive_type: %s\n", primitive_name(info->mesh.primitive_type)); in print_shader_info()
2722 print_nz_bool(fp, "nv", info->mesh.nv); in print_shader_info()
2726 fprintf(fp, "Unhandled stage %d\n", info->stage); in print_shader_info()
2741 print_shader_info(&shader->info, fp); in _nir_print_shader_annotated()
2743 print_nz_unsigned(fp, "inputs", shader->num_inputs); in _nir_print_shader_annotated()
2744 print_nz_unsigned(fp, "outputs", shader->num_outputs); in _nir_print_shader_annotated()
2745 print_nz_unsigned(fp, "uniforms", shader->num_uniforms); in _nir_print_shader_annotated()
2746 if (shader->scratch_size) in _nir_print_shader_annotated()
2747 fprintf(fp, "scratch: %u\n", shader->scratch_size); in _nir_print_shader_annotated()
2748 if (shader->constant_data_size) in _nir_print_shader_annotated()
2749 fprintf(fp, "constants: %u\n", shader->constant_data_size); in _nir_print_shader_annotated()
2759 if (var->data.location == j) in _nir_print_shader_annotated()
2760 vars[var->data.location_frac] = var; in _nir_print_shader_annotated()
2773 foreach_list_typed(nir_function, func, node, &shader->functions) { in _nir_print_shader_annotated()
2835 if (instr->block) { in nir_print_instr()
2836 nir_function_impl *impl = nir_cf_node_get_function(&instr->block->cf_node); in nir_print_instr()
2837 state.shader = impl->function->shader; in nir_print_instr()
2890 instr->index = instr_count; in nir_shader_gather_debug_info()
2903 nir_builder b = nir_builder_at(nir_before_cf_list(&impl->body)); in nir_shader_gather_debug_info()
2908 if (instr->type == nir_instr_type_debug_info || in nir_shader_gather_debug_info()
2909 instr->type == nir_instr_type_phi) in nir_shader_gather_debug_info()
2913 di->src_loc.filename = nir_src_for_ssa(filename_def); in nir_shader_gather_debug_info()
2914 di->src_loc.source = nir_debug_info_nir; in nir_shader_gather_debug_info()
2930 while (character_index < di->src_loc.column) { in nir_shader_gather_debug_info()
2936 di->src_loc.line = line; in nir_shader_gather_debug_info()
2937 di->src_loc.column = 0; in nir_shader_gather_debug_info()
2944 if (instr->type != nir_instr_type_debug_info && in nir_shader_gather_debug_info()
2945 instr->type != nir_instr_type_phi) in nir_shader_gather_debug_info()
2946 nir_instr_insert_before(instr, &debug_info[instr_count++]->instr); in nir_shader_gather_debug_info()