Lines Matching defs:nir
29 #include "nir/nir.h"
30 #include "nir/nir_builder.h"
31 #include "nir/nir_xfb_info.h"
141 is_meta_shader(nir_shader *nir)
143 return nir && nir->info.internal;
147 radv_can_dump_shader(struct radv_device *device, nir_shader *nir, bool meta_shader)
152 if ((is_meta_shader(nir) || meta_shader) &&
160 radv_can_dump_shader_stats(struct radv_device *device, nir_shader *nir)
163 return device->instance->debug_flags & RADV_DEBUG_DUMP_SHADER_STATS && !is_meta_shader(nir);
231 radv_optimize_nir_algebraic(nir_shader *nir, bool opt_offsets)
236 NIR_PASS(_, nir, nir_copy_prop);
237 NIR_PASS(_, nir, nir_opt_dce);
238 NIR_PASS(_, nir, nir_opt_constant_folding);
239 NIR_PASS(_, nir, nir_opt_cse);
240 NIR_PASS(more_algebraic, nir, nir_opt_algebraic);
249 NIR_PASS(_, nir, nir_opt_offsets, &offset_options);
261 NIR_PASS(more_late_algebraic, nir, nir_opt_algebraic_late);
262 NIR_PASS(_, nir, nir_opt_constant_folding);
263 NIR_PASS(_, nir, nir_copy_prop);
264 NIR_PASS(_, nir, nir_opt_dce);
265 NIR_PASS(_, nir, nir_opt_cse);
322 lower_intrinsics(nir_shader *nir, const struct radv_pipeline_key *key)
324 nir_function_impl *entry = nir_shader_get_entrypoint(nir);
371 radv_lower_primitive_shading_rate(nir_shader *nir, enum amd_gfx_level gfx_level)
373 nir_function_impl *impl = nir_shader_get_entrypoint(nir);
431 if (nir->info.stage == MESA_SHADER_MESH) {
441 if (nir->info.stage == MESA_SHADER_VERTEX)
444 if (nir->info.stage == MESA_SHADER_VERTEX && progress)
457 radv_force_primitive_shading_rate(nir_shader *nir, struct radv_device *device)
459 nir_function_impl *impl = nir_shader_get_entrypoint(nir);
489 var = nir_variable_create(nir, nir_var_shader_out, glsl_int_type(), "vrs rate");
503 nir->info.outputs_written |= BITFIELD64_BIT(VARYING_SLOT_PRIMITIVE_SHADING_RATE);
506 if (nir->info.stage == MESA_SHADER_VERTEX)
510 if (nir->info.stage == MESA_SHADER_VERTEX && progress)
523 radv_lower_fs_intrinsics(nir_shader *nir, const struct radv_pipeline_stage *fs_stage,
528 nir_function_impl *impl = nir_shader_get_entrypoint(nir);
614 radv_lower_ms_workgroup_id(nir_shader *nir)
616 nir_function_impl *impl = nir_shader_get_entrypoint(nir);
684 nir_shader *nir;
691 nir = nir_shader_clone(NULL, stage->internal_nir);
692 nir_validate_shader(nir, "in internal shader");
694 assert(exec_list_length(&nir->functions) == 1);
790 nir = spirv_to_nir(spirv, stage->spirv.size / 4, spec_entries, num_spec_entries, stage->stage,
793 nir->info.internal |= device->app_shaders_internal;
794 assert(nir->info.stage == stage->stage);
795 nir_validate_shader(nir, "after spirv_to_nir");
802 NIR_PASS_V(nir, nir_lower_sysvals_to_varyings, &sysvals_to_varyings);
808 NIR_PASS(_, nir, nir_lower_variable_initializers, nir_var_function_temp);
809 NIR_PASS(_, nir, nir_lower_returns);
811 NIR_PASS(progress, nir, nir_inline_functions);
813 NIR_PASS(_, nir, nir_opt_copy_prop_vars);
814 NIR_PASS(_, nir, nir_copy_prop);
816 NIR_PASS(_, nir, nir_opt_deref);
819 foreach_list_typed_safe(nir_function, func, node, &nir->functions)
826 assert(exec_list_length(&nir->functions) == 1);
831 NIR_PASS(_, nir, nir_lower_variable_initializers, nir_var_shader_out);
836 NIR_PASS(_, nir, nir_lower_variable_initializers, ~0);
841 NIR_PASS(_, nir, nir_split_var_copies);
842 NIR_PASS(_, nir, nir_split_per_member_structs);
844 if (nir->info.stage == MESA_SHADER_FRAGMENT)
845 NIR_PASS(_, nir, nir_lower_io_to_vector, nir_var_shader_out);
846 if (nir->info.stage == MESA_SHADER_FRAGMENT)
847 NIR_PASS(_, nir, nir_lower_input_attachments,
856 NIR_PASS(_, nir, nir_remove_dead_variables,
863 NIR_PASS(_, nir, nir_lower_global_vars_to_local);
864 NIR_PASS(_, nir, nir_lower_vars_to_ssa);
866 NIR_PASS(_, nir, nir_propagate_invariant, key->invariant_geom);
868 NIR_PASS(_, nir, nir_lower_clip_cull_distance_arrays);
870 if (nir->info.stage == MESA_SHADER_VERTEX ||
871 nir->info.stage == MESA_SHADER_TESS_EVAL ||
872 nir->info.stage == MESA_SHADER_GEOMETRY)
873 NIR_PASS_V(nir, nir_shader_gather_xfb_info);
875 NIR_PASS(_, nir, nir_lower_discard_or_demote, key->ps.lower_discard_to_demote);
877 nir_lower_doubles_options lower_doubles = nir->options->lower_doubles_options;
887 NIR_PASS(_, nir, nir_lower_doubles, NULL, lower_doubles);
889 NIR_PASS(_, nir, nir_shader_lower_instructions, &is_sincos, &lower_sincos, NULL);
892 NIR_PASS(_, nir, nir_lower_system_values);
897 .lower_cs_local_id_to_index = nir->info.stage == MESA_SHADER_MESH,
898 .lower_local_invocation_index = nir->info.stage == MESA_SHADER_COMPUTE &&
899 ((nir->info.workgroup_size[0] == 1) +
900 (nir->info.workgroup_size[1] == 1) +
901 (nir->info.workgroup_size[2] == 1)) == 2,
903 NIR_PASS(_, nir, nir_lower_compute_system_values, &csv_options);
905 if (nir->info.stage == MESA_SHADER_MESH) {
907 NIR_PASS(_, nir, radv_lower_ms_workgroup_id);
915 NIR_PASS(_, nir, nir_lower_compute_system_values, &o);
919 nir->info.separate_shader = true;
921 nir_shader_gather_info(nir, nir_shader_get_entrypoint(nir));
923 if (nir->info.ray_queries > 0) {
924 NIR_PASS(_, nir, nir_opt_ray_queries);
925 NIR_PASS(_, nir, radv_nir_lower_ray_queries, device);
939 NIR_PASS(_, nir, nir_lower_tex, &tex_options);
945 NIR_PASS(_, nir, nir_lower_image, &image_options);
947 NIR_PASS(_, nir, nir_lower_vars_to_ssa);
949 if (nir->info.stage == MESA_SHADER_VERTEX || nir->info.stage == MESA_SHADER_GEOMETRY ||
950 nir->info.stage == MESA_SHADER_FRAGMENT) {
951 NIR_PASS_V(nir, nir_lower_io_to_temporaries, nir_shader_get_entrypoint(nir), true, true);
952 } else if (nir->info.stage == MESA_SHADER_TESS_EVAL) {
953 NIR_PASS_V(nir, nir_lower_io_to_temporaries, nir_shader_get_entrypoint(nir), true, false);
956 NIR_PASS(_, nir, nir_split_var_copies);
958 NIR_PASS(_, nir, nir_lower_global_vars_to_local);
959 NIR_PASS(_, nir, nir_remove_dead_variables, nir_var_function_temp, NULL);
961 NIR_PASS(_, nir, nir_lower_subgroups,
976 NIR_PASS(_, nir, nir_lower_load_const_to_scalar);
977 NIR_PASS(_, nir, nir_opt_shrink_stores, !device->instance->disable_shrink_image_store);
980 radv_optimize_nir(nir, false, true);
985 NIR_PASS(_, nir, nir_lower_var_copies);
987 unsigned lower_flrp = (nir->options->lower_flrp16 ? 16 : 0) |
988 (nir->options->lower_flrp32 ? 32 : 0) |
989 (nir->options->lower_flrp64 ? 64 : 0);
992 NIR_PASS(progress, nir, nir_lower_flrp, lower_flrp, false /* always precise */);
994 NIR_PASS(_, nir, nir_opt_constant_folding);
1001 NIR_PASS(_, nir, nir_opt_access, &opt_access_options);
1003 NIR_PASS(_, nir, nir_lower_explicit_io, nir_var_mem_push_const, nir_address_format_32bit_offset);
1005 NIR_PASS(_, nir, nir_lower_explicit_io, nir_var_mem_ubo | nir_var_mem_ssbo,
1008 NIR_PASS(_, nir, lower_intrinsics, key);
1011 if (nir->info.stage == MESA_SHADER_COMPUTE ||
1012 nir->info.stage == MESA_SHADER_TASK ||
1013 nir->info.stage == MESA_SHADER_MESH) {
1016 if (nir->info.stage == MESA_SHADER_TASK ||
1017 nir->info.stage == MESA_SHADER_MESH)
1020 if (!nir->info.shared_memory_explicit_layout) {
1021 NIR_PASS(_, nir, nir_lower_vars_to_explicit_types, var_modes, shared_var_info);
1023 NIR_PASS(_, nir, nir_lower_explicit_io, var_modes, nir_address_format_32bit_offset);
1025 if (nir->info.zero_initialize_shared_memory && nir->info.shared_size > 0) {
1027 const unsigned shared_size = ALIGN(nir->info.shared_size, chunk_size);
1028 NIR_PASS(_, nir, nir_zero_initialize_shared_memory, shared_size, chunk_size);
1032 NIR_PASS(_, nir, nir_lower_explicit_io, nir_var_mem_global | nir_var_mem_constant,
1039 NIR_PASS(_, nir, nir_opt_large_constants, glsl_get_natural_size_align_bytes, 16);
1042 if ((nir->info.stage == MESA_SHADER_VERTEX ||
1043 nir->info.stage == MESA_SHADER_GEOMETRY ||
1044 nir->info.stage == MESA_SHADER_MESH) &&
1045 nir->info.outputs_written & BITFIELD64_BIT(VARYING_SLOT_PRIMITIVE_SHADING_RATE)) {
1047 NIR_PASS(_, nir, radv_lower_primitive_shading_rate,
1056 if (ac_nir_lower_indirect_derefs(nir, device->physical_device->rad_info.gfx_level) &&
1057 !key->optimisations_disabled && nir->info.stage != MESA_SHADER_COMPUTE) {
1059 radv_optimize_nir(nir, false, false);
1063 return nir;
1073 find_layer_in_var(nir_shader *nir)
1075 nir_variable *var = nir_find_variable_with_location(nir, nir_var_shader_in, VARYING_SLOT_LAYER);
1079 var = nir_variable_create(nir, nir_var_shader_in, glsl_int_type(), "layer id");
1095 lower_view_index(nir_shader *nir, bool per_primitive)
1098 nir_function_impl *entry = nir_shader_get_entrypoint(nir);
1113 layer = find_layer_in_var(nir);
1121 nir->info.inputs_read |= VARYING_BIT_LAYER;
1123 nir->info.per_primitive_inputs |= VARYING_BIT_LAYER;
1139 radv_lower_io(struct radv_device *device, nir_shader *nir, bool is_mesh_shading)
1141 if (nir->info.stage == MESA_SHADER_COMPUTE)
1144 if (nir->info.stage == MESA_SHADER_FRAGMENT) {
1145 NIR_PASS(_, nir, lower_view_index, is_mesh_shading);
1146 nir_assign_io_var_locations(nir, nir_var_shader_in, &nir->num_inputs, MESA_SHADER_FRAGMENT);
1149 NIR_PASS(_, nir, nir_lower_io, nir_var_shader_in | nir_var_shader_out, type_size_vec4,
1153 NIR_PASS(_, nir, nir_opt_constant_folding);
1155 NIR_PASS(_, nir, nir_io_add_const_offset_to_base, nir_var_shader_in | nir_var_shader_out);
1163 nir_shader *nir = stage->nir;
1165 if (nir->info.stage == MESA_SHADER_VERTEX) {
1167 NIR_PASS_V(nir, ac_nir_lower_ls_outputs_to_mem, NULL, info->vs.tcs_in_out_eq,
1171 NIR_PASS_V(nir, ac_nir_lower_es_outputs_to_mem, NULL,
1176 } else if (nir->info.stage == MESA_SHADER_TESS_CTRL) {
1177 NIR_PASS_V(nir, ac_nir_lower_hs_inputs_to_mem, NULL, info->vs.tcs_in_out_eq);
1178 NIR_PASS_V(nir, ac_nir_lower_hs_outputs_to_mem, NULL,
1186 } else if (nir->info.stage == MESA_SHADER_TESS_EVAL) {
1187 NIR_PASS_V(nir, ac_nir_lower_tes_inputs_to_mem, NULL);
1190 NIR_PASS_V(nir, ac_nir_lower_es_outputs_to_mem, NULL,
1196 } else if (nir->info.stage == MESA_SHADER_GEOMETRY) {
1197 NIR_PASS_V(nir, ac_nir_lower_gs_inputs_to_mem, NULL,
1200 } else if (nir->info.stage == MESA_SHADER_TASK) {
1201 ac_nir_apply_first_task_to_task_shader(nir);
1202 ac_nir_lower_task_outputs_to_mem(nir, AC_TASK_PAYLOAD_ENTRY_BYTES,
1205 } else if (nir->info.stage == MESA_SHADER_MESH) {
1206 ac_nir_lower_mesh_inputs_to_mem(nir, AC_TASK_PAYLOAD_ENTRY_BYTES,
1215 radv_consider_culling(const struct radv_physical_device *pdevice, struct nir_shader *nir, uint64_t ps_inputs_read,
1219 if (is_meta_shader(nir))
1223 if (nir->info.outputs_written & (VARYING_BIT_VIEWPORT | VARYING_BIT_VIEWPORT_MASK))
1259 if (nir->info.writes_memory)
1265 if (BITSET_TEST(nir->info.system_values_read, SYSTEM_VALUE_SUBGROUP_INVOCATION))
1275 nir_shader *nir = ngg_stage->nir;
1277 assert(nir->info.stage == MESA_SHADER_VERTEX ||
1278 nir->info.stage == MESA_SHADER_TESS_EVAL ||
1279 nir->info.stage == MESA_SHADER_GEOMETRY ||
1280 nir->info.stage == MESA_SHADER_MESH);
1286 if (nir->info.stage == MESA_SHADER_TESS_EVAL) {
1287 if (nir->info.tess.point_mode)
1289 else if (nir->info.tess._primitive_mode == TESS_PRIMITIVE_ISOLINES)
1294 BITSET_SET(nir->info.system_values_read, SYSTEM_VALUE_PRIMITIVE_ID);
1296 } else if (nir->info.stage == MESA_SHADER_VERTEX) {
1302 BITSET_SET(nir->info.system_values_read, SYSTEM_VALUE_INSTANCE_ID);
1304 } else if (nir->info.stage == MESA_SHADER_GEOMETRY) {
1305 num_vertices_per_prim = nir->info.gs.vertices_in;
1306 } else if (nir->info.stage == MESA_SHADER_MESH) {
1307 if (nir->info.mesh.primitive_type == SHADER_PRIM_POINTS)
1309 else if (nir->info.mesh.primitive_type == SHADER_PRIM_LINES)
1312 assert(nir->info.mesh.primitive_type == SHADER_PRIM_TRIANGLES);
1320 if (nir->info.stage == MESA_SHADER_VERTEX ||
1321 nir->info.stage == MESA_SHADER_TESS_EVAL) {
1327 radv_optimize_nir_algebraic(nir, false);
1329 if (nir->info.stage == MESA_SHADER_VERTEX) {
1335 NIR_PASS_V(nir, ac_nir_lower_ngg_nogs,
1344 ngg_stage->info.ngg_info.esgs_ring_size = nir->info.shared_size;
1345 } else if (nir->info.stage == MESA_SHADER_GEOMETRY) {
1347 NIR_PASS_V(nir, ac_nir_lower_ngg_gs, info->wave_size, info->workgroup_size,
1350 } else if (nir->info.stage == MESA_SHADER_MESH) {
1352 NIR_PASS_V(nir, ac_nir_lower_ngg_ms, &scratch_ring, info->wave_size, pl_key->has_multiview_view_index);