Lines Matching full:nir
13 #define OPT(nir, pass, ...) ({ \ argument
15 NIR_PASS(this_progress, nir, pass, ##__VA_ARGS__); \
21 #define OPT_V(nir, pass, ...) NIR_PASS_V(nir, pass, ##__VA_ARGS__) argument
24 nak_nir_workgroup_has_one_subgroup(const nir_shader *nir) in nak_nir_workgroup_has_one_subgroup() argument
26 switch (nir->info.stage) { in nak_nir_workgroup_has_one_subgroup()
43 if (nir->info.workgroup_size_variable) in nak_nir_workgroup_has_one_subgroup()
46 uint16_t wg_sz = nir->info.workgroup_size[0] * in nak_nir_workgroup_has_one_subgroup()
47 nir->info.workgroup_size[1] * in nak_nir_workgroup_has_one_subgroup()
48 nir->info.workgroup_size[2]; in nak_nir_workgroup_has_one_subgroup()
92 optimize_nir(nir_shader *nir, const struct nak_compiler *nak, bool allow_copies) in optimize_nir() argument
97 (nir->options->lower_flrp16 ? 16 : 0) | in optimize_nir()
98 (nir->options->lower_flrp32 ? 32 : 0) | in optimize_nir()
99 (nir->options->lower_flrp64 ? 64 : 0); in optimize_nir()
110 if (nir->info.stage != MESA_SHADER_KERNEL) in optimize_nir()
111 OPT(nir, nir_split_array_vars, nir_var_function_temp); in optimize_nir()
113 OPT(nir, nir_shrink_vec_array_vars, nir_var_function_temp); in optimize_nir()
114 OPT(nir, nir_opt_deref); in optimize_nir()
115 if (OPT(nir, nir_opt_memcpy)) in optimize_nir()
116 OPT(nir, nir_split_var_copies); in optimize_nir()
118 OPT(nir, nir_lower_vars_to_ssa); in optimize_nir()
125 OPT(nir, nir_opt_find_array_copies); in optimize_nir()
127 OPT(nir, nir_opt_copy_prop_vars); in optimize_nir()
128 OPT(nir, nir_opt_dead_write_vars); in optimize_nir()
129 OPT(nir, nir_opt_combine_stores, nir_var_all); in optimize_nir()
131 OPT(nir, nir_lower_alu_width, vectorize_filter_cb, NULL); in optimize_nir()
132 OPT(nir, nir_opt_vectorize, vectorize_filter_cb, NULL); in optimize_nir()
133 OPT(nir, nir_lower_phis_to_scalar, false); in optimize_nir()
134 OPT(nir, nir_lower_frexp); in optimize_nir()
135 OPT(nir, nir_copy_prop); in optimize_nir()
136 OPT(nir, nir_opt_dce); in optimize_nir()
137 OPT(nir, nir_opt_cse); in optimize_nir()
139 OPT(nir, nir_opt_peephole_select, 0, false, false); in optimize_nir()
140 OPT(nir, nir_opt_intrinsics); in optimize_nir()
141 OPT(nir, nir_opt_idiv_const, 32); in optimize_nir()
142 OPT(nir, nir_opt_algebraic); in optimize_nir()
143 OPT(nir, nir_lower_constant_convert_alu_types); in optimize_nir()
144 OPT(nir, nir_opt_constant_folding); in optimize_nir()
147 if (OPT(nir, nir_lower_flrp, lower_flrp, false /* always_precise */)) in optimize_nir()
148 OPT(nir, nir_opt_constant_folding); in optimize_nir()
153 OPT(nir, nir_opt_dead_cf); in optimize_nir()
154 if (OPT(nir, nir_opt_loop)) { in optimize_nir()
159 OPT(nir, nir_copy_prop); in optimize_nir()
160 OPT(nir, nir_opt_dce); in optimize_nir()
162 OPT(nir, nir_opt_if, nir_opt_if_optimize_phi_true_false); in optimize_nir()
163 OPT(nir, nir_opt_conditional_discard); in optimize_nir()
164 if (nir->options->max_unroll_iterations != 0) { in optimize_nir()
165 OPT(nir, nir_opt_loop_unroll); in optimize_nir()
167 OPT(nir, nir_opt_remove_phis); in optimize_nir()
168 OPT(nir, nir_opt_gcm, false); in optimize_nir()
169 OPT(nir, nir_opt_undef); in optimize_nir()
170 OPT(nir, nir_lower_pack); in optimize_nir()
173 OPT(nir, nir_remove_dead_variables, nir_var_function_temp, NULL); in optimize_nir()
177 nak_optimize_nir(nir_shader *nir, const struct nak_compiler *nak) in nak_optimize_nir() argument
179 optimize_nir(nir, nak, false); in nak_optimize_nir()
292 nak_preprocess_nir(nir_shader *nir, const struct nak_compiler *nak) in nak_preprocess_nir() argument
296 nir_validate_ssa_dominance(nir, "before nak_preprocess_nir"); in nak_preprocess_nir()
298 if (nir->info.stage == MESA_SHADER_FRAGMENT) { in nak_preprocess_nir()
299 nir_lower_io_to_temporaries(nir, nir_shader_get_entrypoint(nir), in nak_preprocess_nir()
311 OPT(nir, nir_lower_tex, &tex_options); in nak_preprocess_nir()
312 OPT(nir, nir_normalize_cubemap_coords); in nak_preprocess_nir()
317 OPT(nir, nir_lower_image, &image_options); in nak_preprocess_nir()
319 OPT(nir, nir_lower_global_vars_to_local); in nak_preprocess_nir()
321 OPT(nir, nir_split_var_copies); in nak_preprocess_nir()
322 OPT(nir, nir_split_struct_vars, nir_var_function_temp); in nak_preprocess_nir()
325 optimize_nir(nir, nak, true /* allow_copies */); in nak_preprocess_nir()
327 OPT(nir, nir_lower_load_const_to_scalar); in nak_preprocess_nir()
328 OPT(nir, nir_lower_var_copies); in nak_preprocess_nir()
329 OPT(nir, nir_lower_system_values); in nak_preprocess_nir()
330 OPT(nir, nir_lower_compute_system_values, NULL); in nak_preprocess_nir()
332 if (nir->info.stage == MESA_SHADER_FRAGMENT) in nak_preprocess_nir()
333 OPT(nir, nir_lower_terminate_to_demote); in nak_preprocess_nir()
655 nak_nir_lower_system_values(nir_shader *nir, const struct nak_compiler *nak) in nak_nir_lower_system_values() argument
657 return nir_shader_intrinsics_pass(nir, nak_nir_lower_system_value_intrin, in nak_nir_lower_system_values()
727 nak_nir_lower_fs_outputs(nir_shader *nir) in nak_nir_lower_fs_outputs() argument
729 if (nir->info.outputs_written == 0) in nak_nir_lower_fs_outputs()
732 bool progress = nir_shader_intrinsics_pass(nir, lower_fs_output_intrin, in nak_nir_lower_fs_outputs()
741 nir_function_impl *impl = nir_shader_get_entrypoint(nir); in nak_nir_lower_fs_outputs()
789 nak_nir_remove_barriers(nir_shader *nir) in nak_nir_remove_barriers() argument
792 nir->info.uses_control_barrier = false; in nak_nir_remove_barriers()
794 return nir_shader_intrinsics_pass(nir, nak_nir_remove_barrier_intrin, in nak_nir_remove_barriers()
890 nir_shader_has_local_variables(const nir_shader *nir) in nir_shader_has_local_variables() argument
892 nir_foreach_function(func, nir) { in nir_shader_has_local_variables()
907 nak_postprocess_nir(nir_shader *nir, in nak_postprocess_nir() argument
914 nak_optimize_nir(nir, nak); in nak_postprocess_nir()
929 OPT(nir, nir_lower_subgroups, &subgroups_options); in nak_postprocess_nir()
930 OPT(nir, nak_nir_lower_scan_reduce); in nak_postprocess_nir()
932 if (nir_shader_has_local_variables(nir)) { in nak_postprocess_nir()
933 OPT(nir, nir_lower_vars_to_explicit_types, nir_var_function_temp, in nak_postprocess_nir()
935 OPT(nir, nir_lower_explicit_io, nir_var_function_temp, in nak_postprocess_nir()
937 nak_optimize_nir(nir, nak); in nak_postprocess_nir()
940 OPT(nir, nir_opt_shrink_vectors, true); in nak_postprocess_nir()
949 OPT(nir, nir_opt_load_store_vectorize, &vectorize_opts); in nak_postprocess_nir()
955 OPT(nir, nir_lower_mem_access_bit_sizes, &mem_bit_size_options); in nak_postprocess_nir()
956 OPT(nir, nir_lower_bit_size, lower_bit_size_cb, (void *)nak); in nak_postprocess_nir()
958 OPT(nir, nir_opt_combine_barriers, NULL, NULL); in nak_postprocess_nir()
960 nak_optimize_nir(nir, nak); in nak_postprocess_nir()
962 OPT(nir, nak_nir_lower_tex, nak); in nak_postprocess_nir()
963 OPT(nir, nir_lower_idiv, NULL); in nak_postprocess_nir()
965 nir_shader_gather_info(nir, nir_shader_get_entrypoint(nir)); in nak_postprocess_nir()
967 OPT(nir, nir_lower_indirect_derefs, 0, UINT32_MAX); in nak_postprocess_nir()
969 if (nir->info.stage == MESA_SHADER_TESS_EVAL) { in nak_postprocess_nir()
970 OPT(nir, nir_lower_tess_coord_z, in nak_postprocess_nir()
971 nir->info.tess._primitive_mode == TESS_PRIMITIVE_TRIANGLES); in nak_postprocess_nir()
978 if (gl_shader_stage_uses_workgroup(nir->info.stage) && in nak_postprocess_nir()
979 nir->info.derivative_group == DERIVATIVE_GROUP_QUADS) { in nak_postprocess_nir()
980 assert(nir->info.workgroup_size[0] % 2 == 0); in nak_postprocess_nir()
981 assert(nir->info.workgroup_size[1] % 2 == 0); in nak_postprocess_nir()
982 nir->info.workgroup_size[0] *= 2; in nak_postprocess_nir()
983 nir->info.workgroup_size[1] /= 2; in nak_postprocess_nir()
986 OPT(nir, nak_nir_lower_system_values, nak); in nak_postprocess_nir()
988 switch (nir->info.stage) { in nak_postprocess_nir()
993 OPT(nir, nir_lower_io, nir_var_shader_in | nir_var_shader_out, in nak_postprocess_nir()
995 OPT(nir, nir_opt_constant_folding); in nak_postprocess_nir()
996 OPT(nir, nak_nir_lower_vtg_io, nak); in nak_postprocess_nir()
997 if (nir->info.stage == MESA_SHADER_GEOMETRY) in nak_postprocess_nir()
998 OPT(nir, nak_nir_lower_gs_intrinsics); in nak_postprocess_nir()
1002 OPT(nir, nir_lower_indirect_derefs, in nak_postprocess_nir()
1004 OPT(nir, nir_lower_io, nir_var_shader_in | nir_var_shader_out, in nak_postprocess_nir()
1007 OPT(nir, nir_opt_constant_folding); in nak_postprocess_nir()
1008 OPT(nir, nak_nir_lower_fs_inputs, nak, fs_key); in nak_postprocess_nir()
1009 OPT(nir, nak_nir_lower_fs_outputs); in nak_postprocess_nir()
1020 OPT(nir, nir_lower_doubles, NULL, nak->nir_options.lower_doubles_options); in nak_postprocess_nir()
1021 OPT(nir, nir_lower_int64); in nak_postprocess_nir()
1023 nak_optimize_nir(nir, nak); in nak_postprocess_nir()
1027 OPT(nir, nir_opt_algebraic_late); in nak_postprocess_nir()
1028 OPT(nir, nak_nir_lower_algebraic_late, nak); in nak_postprocess_nir()
1035 OPT(nir, nir_lower_doubles, NULL, nir_lower_dsat); in nak_postprocess_nir()
1038 OPT(nir, nir_opt_constant_folding); in nak_postprocess_nir()
1039 OPT(nir, nir_copy_prop); in nak_postprocess_nir()
1040 OPT(nir, nir_opt_dce); in nak_postprocess_nir()
1041 OPT(nir, nir_opt_cse); in nak_postprocess_nir()
1046 OPT(nir, nak_nir_split_64bit_conversions); in nak_postprocess_nir()
1048 bool lcssa_progress = nir_convert_to_lcssa(nir, false, false); in nak_postprocess_nir()
1049 nir_divergence_analysis(nir); in nak_postprocess_nir()
1053 OPT(nir, nak_nir_mark_lcssa_invariants); in nak_postprocess_nir()
1055 if (OPT(nir, nak_nir_lower_non_uniform_ldcx)) { in nak_postprocess_nir()
1056 OPT(nir, nir_copy_prop); in nak_postprocess_nir()
1057 OPT(nir, nir_opt_dce); in nak_postprocess_nir()
1058 nir_divergence_analysis(nir); in nak_postprocess_nir()
1062 OPT(nir, nak_nir_remove_barriers); in nak_postprocess_nir()
1066 fprintf(stderr, "Structured NIR for %s shader:\n", in nak_postprocess_nir()
1067 _mesa_shader_stage_to_string(nir->info.stage)); in nak_postprocess_nir()
1068 nir_print_shader(nir, stderr); in nak_postprocess_nir()
1070 OPT(nir, nak_nir_lower_cf); in nak_postprocess_nir()
1076 nir_foreach_function(func, nir) { in nak_postprocess_nir()
1084 fprintf(stderr, "NIR for %s shader:\n", in nak_postprocess_nir()
1085 _mesa_shader_stage_to_string(nir->info.stage)); in nak_postprocess_nir()
1086 nir_print_shader(nir, stderr); in nak_postprocess_nir()