Lines Matching full:nir
12 #define OPT(nir, pass, ...) ({ \ argument
14 NIR_PASS(this_progress, nir, pass, ##__VA_ARGS__); \
20 #define OPT_V(nir, pass, ...) NIR_PASS_V(nir, pass, ##__VA_ARGS__) argument
23 nak_nir_workgroup_has_one_subgroup(const nir_shader *nir) in nak_nir_workgroup_has_one_subgroup() argument
25 switch (nir->info.stage) { in nak_nir_workgroup_has_one_subgroup()
42 if (nir->info.workgroup_size_variable) in nak_nir_workgroup_has_one_subgroup()
45 uint16_t wg_sz = nir->info.workgroup_size[0] * in nak_nir_workgroup_has_one_subgroup()
46 nir->info.workgroup_size[1] * in nak_nir_workgroup_has_one_subgroup()
47 nir->info.workgroup_size[2]; in nak_nir_workgroup_has_one_subgroup()
58 optimize_nir(nir_shader *nir, const struct nak_compiler *nak, bool allow_copies) in optimize_nir() argument
63 (nir->options->lower_flrp16 ? 16 : 0) | in optimize_nir()
64 (nir->options->lower_flrp32 ? 32 : 0) | in optimize_nir()
65 (nir->options->lower_flrp64 ? 64 : 0); in optimize_nir()
76 if (nir->info.stage != MESA_SHADER_KERNEL) in optimize_nir()
77 OPT(nir, nir_split_array_vars, nir_var_function_temp); in optimize_nir()
79 OPT(nir, nir_shrink_vec_array_vars, nir_var_function_temp); in optimize_nir()
80 OPT(nir, nir_opt_deref); in optimize_nir()
81 if (OPT(nir, nir_opt_memcpy)) in optimize_nir()
82 OPT(nir, nir_split_var_copies); in optimize_nir()
84 OPT(nir, nir_lower_vars_to_ssa); in optimize_nir()
91 OPT(nir, nir_opt_find_array_copies); in optimize_nir()
93 OPT(nir, nir_opt_copy_prop_vars); in optimize_nir()
94 OPT(nir, nir_opt_dead_write_vars); in optimize_nir()
95 OPT(nir, nir_opt_combine_stores, nir_var_all); in optimize_nir()
97 OPT(nir, nir_lower_alu_to_scalar, NULL, NULL); in optimize_nir()
98 OPT(nir, nir_lower_phis_to_scalar, false); in optimize_nir()
99 OPT(nir, nir_lower_frexp); in optimize_nir()
100 OPT(nir, nir_copy_prop); in optimize_nir()
101 OPT(nir, nir_opt_dce); in optimize_nir()
102 OPT(nir, nir_opt_cse); in optimize_nir()
104 OPT(nir, nir_opt_peephole_select, 0, false, false); in optimize_nir()
105 OPT(nir, nir_opt_intrinsics); in optimize_nir()
106 OPT(nir, nir_opt_idiv_const, 32); in optimize_nir()
107 OPT(nir, nir_opt_algebraic); in optimize_nir()
108 OPT(nir, nir_lower_constant_convert_alu_types); in optimize_nir()
109 OPT(nir, nir_opt_constant_folding); in optimize_nir()
112 if (OPT(nir, nir_lower_flrp, lower_flrp, false /* always_precise */)) in optimize_nir()
113 OPT(nir, nir_opt_constant_folding); in optimize_nir()
118 OPT(nir, nir_opt_dead_cf); in optimize_nir()
119 if (OPT(nir, nir_opt_loop)) { in optimize_nir()
124 OPT(nir, nir_copy_prop); in optimize_nir()
125 OPT(nir, nir_opt_dce); in optimize_nir()
127 OPT(nir, nir_opt_if, nir_opt_if_optimize_phi_true_false); in optimize_nir()
128 OPT(nir, nir_opt_conditional_discard); in optimize_nir()
129 if (nir->options->max_unroll_iterations != 0) { in optimize_nir()
130 OPT(nir, nir_opt_loop_unroll); in optimize_nir()
132 OPT(nir, nir_opt_remove_phis); in optimize_nir()
133 OPT(nir, nir_opt_gcm, false); in optimize_nir()
134 OPT(nir, nir_opt_undef); in optimize_nir()
135 OPT(nir, nir_lower_pack); in optimize_nir()
138 OPT(nir, nir_remove_dead_variables, nir_var_function_temp, NULL); in optimize_nir()
142 nak_optimize_nir(nir_shader *nir, const struct nak_compiler *nak) in nak_optimize_nir() argument
144 optimize_nir(nir, nak, false); in nak_optimize_nir()
282 nak_nir_lower_subgroup_id(nir_shader *nir) in nak_nir_lower_subgroup_id() argument
284 return nir_shader_intrinsics_pass(nir, nak_nir_lower_subgroup_id_intrin, in nak_nir_lower_subgroup_id()
291 nak_preprocess_nir(nir_shader *nir, const struct nak_compiler *nak) in nak_preprocess_nir() argument
295 nir_validate_ssa_dominance(nir, "before nak_preprocess_nir"); in nak_preprocess_nir()
305 OPT(nir, nir_lower_tex, &tex_options); in nak_preprocess_nir()
306 OPT(nir, nir_normalize_cubemap_coords); in nak_preprocess_nir()
311 OPT(nir, nir_lower_image, &image_options); in nak_preprocess_nir()
313 OPT(nir, nir_lower_global_vars_to_local); in nak_preprocess_nir()
315 OPT(nir, nir_split_var_copies); in nak_preprocess_nir()
316 OPT(nir, nir_split_struct_vars, nir_var_function_temp); in nak_preprocess_nir()
319 optimize_nir(nir, nak, true /* allow_copies */); in nak_preprocess_nir()
321 OPT(nir, nir_lower_load_const_to_scalar); in nak_preprocess_nir()
322 OPT(nir, nir_lower_var_copies); in nak_preprocess_nir()
323 OPT(nir, nir_lower_system_values); in nak_preprocess_nir()
324 OPT(nir, nak_nir_lower_subgroup_id); in nak_preprocess_nir()
325 OPT(nir, nir_lower_compute_system_values, NULL); in nak_preprocess_nir()
342 nak_nir_lower_vs_inputs(nir_shader *nir) in nak_nir_lower_vs_inputs() argument
346 nir_foreach_shader_in_variable(var, nir) { in nak_nir_lower_vs_inputs()
351 progress |= OPT(nir, nir_lower_io, nir_var_shader_in, type_size_vec4_bytes, in nak_nir_lower_vs_inputs()
597 nak_nir_lower_system_values(nir_shader *nir, const struct nak_compiler *nak) in nak_nir_lower_system_values() argument
599 return nir_shader_intrinsics_pass(nir, nak_nir_lower_system_value_intrin, in nak_nir_lower_system_values()
605 nak_nir_lower_varyings(nir_shader *nir, nir_variable_mode modes) in nak_nir_lower_varyings() argument
611 nir_foreach_variable_with_modes(var, nir, modes) in nak_nir_lower_varyings()
614 OPT(nir, nir_lower_io, modes, type_size_vec4_bytes, in nak_nir_lower_varyings()
1010 nak_nir_lower_fs_inputs(nir_shader *nir, in nak_nir_lower_fs_inputs() argument
1014 NIR_PASS_V(nir, nir_lower_indirect_derefs, nir_var_shader_in, UINT32_MAX); in nak_nir_lower_fs_inputs()
1015 NIR_PASS_V(nir, nak_nir_lower_varyings, nir_var_shader_in); in nak_nir_lower_fs_inputs()
1016 NIR_PASS_V(nir, nir_opt_constant_folding); in nak_nir_lower_fs_inputs()
1022 NIR_PASS_V(nir, nir_shader_intrinsics_pass, lower_fs_input_intrin, in nak_nir_lower_fs_inputs()
1037 nak_nir_lower_fs_outputs(nir_shader *nir) in nak_nir_lower_fs_outputs() argument
1039 if (nir->info.outputs_written == 0) in nak_nir_lower_fs_outputs()
1042 NIR_PASS_V(nir, nir_lower_io_arrays_to_elements_no_indirects, true); in nak_nir_lower_fs_outputs()
1044 nir->num_outputs = 0; in nak_nir_lower_fs_outputs()
1045 nir_foreach_shader_out_variable(var, nir) { in nak_nir_lower_fs_outputs()
1074 NIR_PASS_V(nir, nir_lower_io, nir_var_shader_out, fs_out_size, 0); in nak_nir_lower_fs_outputs()
1159 nir_shader_has_local_variables(const nir_shader *nir) in nir_shader_has_local_variables() argument
1161 nir_foreach_function(func, nir) { in nir_shader_has_local_variables()
1170 nak_postprocess_nir(nir_shader *nir, in nak_postprocess_nir() argument
1177 nak_optimize_nir(nir, nak); in nak_postprocess_nir()
1190 OPT(nir, nir_lower_subgroups, &subgroups_options); in nak_postprocess_nir()
1191 OPT(nir, nak_nir_lower_scan_reduce); in nak_postprocess_nir()
1193 if (nir_shader_has_local_variables(nir)) { in nak_postprocess_nir()
1194 OPT(nir, nir_lower_vars_to_explicit_types, nir_var_function_temp, in nak_postprocess_nir()
1196 OPT(nir, nir_lower_explicit_io, nir_var_function_temp, in nak_postprocess_nir()
1198 nak_optimize_nir(nir, nak); in nak_postprocess_nir()
1201 OPT(nir, nir_opt_shrink_vectors); in nak_postprocess_nir()
1210 OPT(nir, nir_opt_load_store_vectorize, &vectorize_opts); in nak_postprocess_nir()
1216 OPT(nir, nir_lower_mem_access_bit_sizes, &mem_bit_size_options); in nak_postprocess_nir()
1217 OPT(nir, nir_lower_bit_size, lower_bit_size_cb, (void *)nak); in nak_postprocess_nir()
1219 OPT(nir, nir_opt_combine_barriers, NULL, NULL); in nak_postprocess_nir()
1221 nak_optimize_nir(nir, nak); in nak_postprocess_nir()
1223 OPT(nir, nak_nir_lower_tex, nak); in nak_postprocess_nir()
1224 OPT(nir, nir_lower_idiv, NULL); in nak_postprocess_nir()
1226 nir_shader_gather_info(nir, nir_shader_get_entrypoint(nir)); in nak_postprocess_nir()
1228 OPT(nir, nir_lower_indirect_derefs, 0, UINT32_MAX); in nak_postprocess_nir()
1230 if (nir->info.stage == MESA_SHADER_TESS_EVAL) { in nak_postprocess_nir()
1231 OPT(nir, nir_lower_tess_coord_z, in nak_postprocess_nir()
1232 nir->info.tess._primitive_mode == TESS_PRIMITIVE_TRIANGLES); in nak_postprocess_nir()
1235 OPT(nir, nak_nir_lower_system_values, nak); in nak_postprocess_nir()
1237 switch (nir->info.stage) { in nak_postprocess_nir()
1239 OPT(nir, nak_nir_lower_vs_inputs); in nak_postprocess_nir()
1240 OPT(nir, nak_nir_lower_varyings, nir_var_shader_out); in nak_postprocess_nir()
1241 OPT(nir, nir_opt_constant_folding); in nak_postprocess_nir()
1242 OPT(nir, nak_nir_lower_vtg_io, nak); in nak_postprocess_nir()
1247 OPT(nir, nak_nir_lower_varyings, nir_var_shader_in | nir_var_shader_out); in nak_postprocess_nir()
1248 OPT(nir, nir_opt_constant_folding); in nak_postprocess_nir()
1249 OPT(nir, nak_nir_lower_vtg_io, nak); in nak_postprocess_nir()
1253 OPT(nir, nak_nir_lower_fs_inputs, nak, fs_key); in nak_postprocess_nir()
1254 OPT(nir, nak_nir_lower_fs_outputs); in nak_postprocess_nir()
1258 OPT(nir, nak_nir_lower_varyings, nir_var_shader_in | nir_var_shader_out); in nak_postprocess_nir()
1259 OPT(nir, nir_opt_constant_folding); in nak_postprocess_nir()
1260 OPT(nir, nak_nir_lower_vtg_io, nak); in nak_postprocess_nir()
1261 OPT(nir, nak_nir_lower_gs_intrinsics); in nak_postprocess_nir()
1272 OPT(nir, nir_lower_doubles, NULL, nak->nir_options.lower_doubles_options); in nak_postprocess_nir()
1273 OPT(nir, nir_lower_int64); in nak_postprocess_nir()
1275 nak_optimize_nir(nir, nak); in nak_postprocess_nir()
1279 OPT(nir, nir_opt_algebraic_late); in nak_postprocess_nir()
1280 OPT(nir, nak_nir_lower_algebraic_late, nak); in nak_postprocess_nir()
1287 OPT(nir, nir_lower_doubles, NULL, nir_lower_dsat); in nak_postprocess_nir()
1290 OPT(nir, nir_opt_constant_folding); in nak_postprocess_nir()
1291 OPT(nir, nir_copy_prop); in nak_postprocess_nir()
1292 OPT(nir, nir_opt_dce); in nak_postprocess_nir()
1293 OPT(nir, nir_opt_cse); in nak_postprocess_nir()
1297 nir_divergence_analysis(nir); in nak_postprocess_nir()
1299 OPT(nir, nak_nir_add_barriers, nak); in nak_postprocess_nir()
1304 nir_foreach_function(func, nir) { in nak_postprocess_nir()
1312 nir_print_shader(nir, stderr); in nak_postprocess_nir()