diff --git a/src/imagination/pco/pco_nir.c b/src/imagination/pco/pco_nir.c index 64d5054045e..275b6bfc3a4 100644 --- a/src/imagination/pco/pco_nir.c +++ b/src/imagination/pco/pco_nir.c @@ -141,7 +141,8 @@ static uint8_t vectorize_filter(const nir_instr *instr, UNUSED const void *data) * \param[in] data User data. * \return True if the instruction was found. */ -static bool frag_in_scalar_filter(const nir_intrinsic_instr *intr, const void *data) +static bool frag_in_scalar_filter(const nir_intrinsic_instr *intr, + const void *data) { nir_shader *nir = (nir_shader *)data; @@ -629,9 +630,7 @@ void pco_preprocess_nir(pco_ctx *ctx, nir_shader *nir) NULL); /* Fold constant offset srcs for IO. */ - NIR_PASS(_, - nir, - nir_opt_constant_folding); + NIR_PASS(_, nir, nir_opt_constant_folding); NIR_PASS(_, nir, @@ -1001,36 +1000,20 @@ void pco_lower_nir(pco_ctx *ctx, nir_shader *nir, pco_data *data) } } -static bool is_phi_with_undefs(const nir_instr *instr, - UNUSED const void *cb_data) +static bool +lower_phi_with_undefs(nir_builder *b, nir_phi_instr *phi, UNUSED void *cb_data) { - if (instr->type != nir_instr_type_phi) - return false; - - nir_phi_instr *phi = nir_instr_as_phi(instr); - - nir_foreach_phi_src (phi_src, phi) { - if (nir_src_is_undef(phi_src->src)) - return true; - } - - return false; -} - -static nir_def * -lower_phi_with_undefs(nir_builder *b, nir_instr *instr, UNUSED void *cb_data) -{ - nir_phi_instr *phi = nir_instr_as_phi(instr); - + bool progress = false; nir_foreach_phi_src (phi_src, phi) { if (nir_src_is_undef(phi_src->src)) { b->cursor = nir_after_block(phi_src->pred); nir_src_rewrite(&phi_src->src, nir_imm_intN_t(b, 0, phi_src->src.ssa->bit_size)); + progress = true; } } - return NIR_LOWER_INSTR_PROGRESS; + return progress; } static bool @@ -1094,10 +1077,10 @@ void pco_postprocess_nir(pco_ctx *ctx, nir_shader *nir, pco_data *data) /* Temporary: lower phi undefs to zero because at this stage we don't want to * lower *all* undefs to zero, but still want to avoid undefined behaviour... */ - nir_shader_lower_instructions(nir, - is_phi_with_undefs, - lower_phi_with_undefs, - NULL); + nir_shader_phi_pass(nir, + lower_phi_with_undefs, + nir_metadata_control_flow, + NULL); NIR_PASS(_, nir, nir_convert_from_ssa, true, false); NIR_PASS(_, nir, nir_opt_copy_prop); diff --git a/src/imagination/pco/pco_nir_io.c b/src/imagination/pco/pco_nir_io.c index 4ab68f3e51b..ea480f3ae19 100644 --- a/src/imagination/pco/pco_nir_io.c +++ b/src/imagination/pco/pco_nir_io.c @@ -25,17 +25,14 @@ * \brief Lowers an I/O instruction. * * \param[in] b NIR builder. - * \param[in] instr NIR instruction. + * \param[in] intr NIR intrinsic instruction. * \param[in] cb_data User callback data. - * \return The replacement/lowered def. + * \return True if progress was made. */ -static nir_def *lower_io(nir_builder *b, nir_instr *instr, UNUSED void *cb_data) +static bool +lower_io(nir_builder *b, nir_intrinsic_instr *intr, UNUSED void *cb_data) { - nir_intrinsic_instr *intr = nir_instr_as_intrinsic(instr); - b->cursor = nir_before_instr(instr); - - ASSERTED unsigned base = nir_intrinsic_base(intr); - assert(!base); + b->cursor = nir_before_instr(&intr->instr); nir_src *offset_src; switch (intr->intrinsic) { @@ -51,41 +48,16 @@ static nir_def *lower_io(nir_builder *b, nir_instr *instr, UNUSED void *cb_data) break; default: - UNREACHABLE(""); + return false; } + ASSERTED unsigned base = nir_intrinsic_base(intr); + assert(!base); + /* Byte offset to DWORD offset. */ nir_src_rewrite(offset_src, nir_ushr_imm(b, offset_src->ssa, 2)); - return NIR_LOWER_INSTR_PROGRESS; -} - -/** - * \brief Filters I/O instructions that need lowering. - * - * \param[in] instr NIR instruction. - * \param[in] cb_data User callback data. - * \return True if the instruction matches the filter. - */ -static bool is_lowerable_io(const nir_instr *instr, UNUSED const void *cb_data) -{ - if (instr->type != nir_instr_type_intrinsic) - return false; - - nir_intrinsic_instr *intr = nir_instr_as_intrinsic(instr); - switch (intr->intrinsic) { - case nir_intrinsic_load_push_constant: - case nir_intrinsic_load_shared: - case nir_intrinsic_store_shared: - case nir_intrinsic_shared_atomic: - case nir_intrinsic_shared_atomic_swap: - return true; - - default: - break; - } - - return false; + return true; } /** @@ -98,8 +70,10 @@ bool pco_nir_lower_io(nir_shader *shader) { bool progress = false; - progress |= - nir_shader_lower_instructions(shader, is_lowerable_io, lower_io, NULL); + progress |= nir_shader_intrinsics_pass(shader, + lower_io, + nir_metadata_control_flow, + NULL); return progress; } diff --git a/src/imagination/pco/pco_nir_pvfio.c b/src/imagination/pco/pco_nir_pvfio.c index fb206d98791..5541f0d2851 100644 --- a/src/imagination/pco/pco_nir_pvfio.c +++ b/src/imagination/pco/pco_nir_pvfio.c @@ -315,14 +315,14 @@ static inline bool is_processed(nir_intrinsic_instr *intr) return nir_alu_type_get_base_type(type) == nir_type_invalid; } -static nir_def *lower_pfo_store(nir_builder *b, - nir_intrinsic_instr *intr, - struct pfo_state *state) +static bool lower_pfo_store(nir_builder *b, + nir_intrinsic_instr *intr, + struct pfo_state *state) { /* Skip stores we've already processed. */ if (is_processed(intr)) { util_dynarray_append(&state->stores, intr); - return NULL; + return false; } nir_def *input = intr->src[0].ssa; @@ -341,8 +341,10 @@ static nir_def *lower_pfo_store(nir_builder *b, b->cursor = nir_before_instr(&intr->instr); enum pipe_format format = state->fs->output_formats[location]; - if (format == PIPE_FORMAT_NONE) - return NIR_LOWER_INSTR_PROGRESS_REPLACE; + if (format == PIPE_FORMAT_NONE) { + nir_instr_remove(&intr->instr); + return true; + } format = to_pbe_format(b, format, &input); @@ -369,17 +371,18 @@ static nir_def *lower_pfo_store(nir_builder *b, assert(var); var->type = glsl_uvec_type(output->num_components); - return NIR_LOWER_INSTR_PROGRESS_REPLACE; + nir_instr_remove(&intr->instr); + return true; } -static nir_def *lower_pfo_load(nir_builder *b, - nir_intrinsic_instr *intr, - struct pfo_state *state) +static bool lower_pfo_load(nir_builder *b, + nir_intrinsic_instr *intr, + struct pfo_state *state) { /* Skip loads we've already processed. */ if (is_processed(intr)) { util_dynarray_append(&state->loads, intr); - return NULL; + return false; } unsigned base = nir_intrinsic_base(intr); @@ -400,75 +403,55 @@ static nir_def *lower_pfo_load(nir_builder *b, format = state->fs->output_formats[location]; } + nir_def *repl; if (format == PIPE_FORMAT_NONE) - return nir_undef(b, intr->def.num_components, intr->def.bit_size); + repl = nir_undef(b, intr->def.num_components, intr->def.bit_size); + else { + format = to_pbe_format(b, format, NULL); - format = to_pbe_format(b, format, NULL); + nir_def *packed_comps[4]; + for (unsigned c = 0; c < ARRAY_SIZE(packed_comps); ++c) { + packed_comps[c] = nir_load_output(b, + 1, + 32, + offset->ssa, + .base = base, + .component = c, + .dest_type = nir_type_invalid | 32, + .io_semantics = io_semantics); - nir_def *packed_comps[4]; - for (unsigned c = 0; c < ARRAY_SIZE(packed_comps); ++c) { - packed_comps[c] = nir_load_output(b, - 1, - 32, - offset->ssa, - .base = base, - .component = c, - .dest_type = nir_type_invalid | 32, - .io_semantics = io_semantics); + nir_intrinsic_instr *load = nir_def_as_intrinsic(packed_comps[c]); - nir_intrinsic_instr *load = - nir_def_as_intrinsic(packed_comps[c]); + util_dynarray_append(&state->loads, load); + } - util_dynarray_append(&state->loads, load); + nir_alu_type dest_type = nir_intrinsic_dest_type(intr); + repl = unpack_from_format(b, + packed_comps, + dest_type, + format, + intr->def.num_components); } - nir_alu_type dest_type = nir_intrinsic_dest_type(intr); - return unpack_from_format(b, - packed_comps, - dest_type, - format, - intr->def.num_components); -} + nir_def_rewrite_uses(&intr->def, repl); + nir_instr_remove(&intr->instr); -/** - * \brief Filters PFO-related instructions. - * - * \param[in] instr NIR instruction. - * \param[in] cb_data User callback data. - * \return True if the instruction matches the filter. - */ -static bool is_pfo(const nir_instr *instr, UNUSED const void *cb_data) -{ - if (instr->type != nir_instr_type_intrinsic) - return false; - - nir_intrinsic_instr *intr = nir_instr_as_intrinsic(instr); - switch (intr->intrinsic) { - case nir_intrinsic_store_output: - case nir_intrinsic_load_output: - case nir_intrinsic_demote: - case nir_intrinsic_demote_if: - return true; - - default: - break; - } - - return false; + return true; } /** * \brief Lowers a PFO-related instruction. * * \param[in] b NIR builder. - * \param[in] instr NIR instruction. + * \param[in] intr NIR intrinsic instruction. * \param[in] cb_data User callback data. - * \return The replacement/lowered def. + * \return True if progress was made. */ -static nir_def *lower_pfo(nir_builder *b, nir_instr *instr, void *cb_data) +static bool lower_pfo(nir_builder *b, nir_intrinsic_instr *intr, void *cb_data) { struct pfo_state *state = cb_data; - nir_intrinsic_instr *intr = nir_instr_as_intrinsic(instr); + + b->cursor = nir_before_instr(&intr->instr); switch (intr->intrinsic) { case nir_intrinsic_store_output: { @@ -480,7 +463,8 @@ static nir_def *lower_pfo(nir_builder *b, nir_instr *instr, void *cb_data) assert(!state->depth_feedback_src); state->depth_feedback_src = nir_fsat(b, intr->src[0].ssa); - return NIR_LOWER_INSTR_PROGRESS_REPLACE; + nir_instr_remove(&intr->instr); + return true; } if (sem.location == FRAG_RESULT_SAMPLE_MASK) { @@ -497,7 +481,8 @@ static nir_def *lower_pfo(nir_builder *b, nir_instr *instr, void *cb_data) val = nir_ior(b, val, cond); state->last_discard_store = nir_build_store_reg(b, val, state->discard_cond_reg); - return NIR_LOWER_INSTR_PROGRESS_REPLACE; + nir_instr_remove(&intr->instr); + return true; } UNREACHABLE(""); @@ -510,7 +495,8 @@ static nir_def *lower_pfo(nir_builder *b, nir_instr *instr, void *cb_data) state->has_discards = true; state->last_discard_store = nir_build_store_reg(b, nir_imm_true(b), state->discard_cond_reg); - return NIR_LOWER_INSTR_PROGRESS_REPLACE; + nir_instr_remove(&intr->instr); + return true; case nir_intrinsic_demote_if: { state->has_discards = true; @@ -518,7 +504,8 @@ static nir_def *lower_pfo(nir_builder *b, nir_instr *instr, void *cb_data) val = nir_ior(b, val, intr->src[0].ssa); state->last_discard_store = nir_build_store_reg(b, val, state->discard_cond_reg); - return NIR_LOWER_INSTR_PROGRESS_REPLACE; + nir_instr_remove(&intr->instr); + return true; } default: @@ -668,20 +655,6 @@ static bool z_replicate(nir_shader *shader, struct pfo_state *state) return true; } -static bool is_frag_color_out(const nir_instr *instr, - UNUSED const void *cb_data) -{ - if (instr->type != nir_instr_type_intrinsic) - return false; - - nir_intrinsic_instr *intr = nir_instr_as_intrinsic(instr); - if (intr->intrinsic != nir_intrinsic_store_output) - return false; - - gl_frag_result location = nir_intrinsic_io_semantics(intr).location; - return location >= FRAG_RESULT_DATA0 && location < FRAG_RESULT_MAX; -} - static bool lower_demote_samples(nir_builder *b, nir_intrinsic_instr *intr, UNUSED void *cb_data) @@ -725,10 +698,16 @@ bool pco_nir_lower_alpha_to_coverage(nir_shader *shader) return true; } -static nir_def * -lower_alpha_to_one(nir_builder *b, nir_instr *instr, UNUSED void *cb_data) +static bool lower_alpha_to_one(nir_builder *b, + nir_intrinsic_instr *intr, + UNUSED void *cb_data) { - nir_intrinsic_instr *intr = nir_instr_as_intrinsic(instr); + if (intr->intrinsic != nir_intrinsic_store_output) + return false; + + gl_frag_result location = nir_intrinsic_io_semantics(intr).location; + if (!(location >= FRAG_RESULT_DATA0 && location < FRAG_RESULT_MAX)) + return false; nir_src *input_src = &intr->src[0]; nir_def *input = input_src->ssa; @@ -737,7 +716,7 @@ lower_alpha_to_one(nir_builder *b, nir_instr *instr, UNUSED void *cb_data) /* Skip color write that don't include alpha. */ if (input->num_components != 4) - return NULL; + return false; b->cursor = nir_before_instr(&intr->instr); @@ -754,7 +733,7 @@ lower_alpha_to_one(nir_builder *b, nir_instr *instr, UNUSED void *cb_data) nir_src_rewrite(input_src, nir_vector_insert_imm(b, input, alpha, 3)); - return NIR_LOWER_INSTR_PROGRESS; + return true; } static bool is_load_sample_mask(const nir_instr *instr, @@ -767,23 +746,33 @@ static bool is_load_sample_mask(const nir_instr *instr, return intr->intrinsic == nir_intrinsic_load_sample_mask_in; } -static nir_def * -lower_load_sample_mask(nir_builder *b, nir_instr *instr, UNUSED void *cb_data) +static bool lower_load_sample_mask(nir_builder *b, + nir_intrinsic_instr *intr, + UNUSED void *cb_data) { - b->cursor = nir_before_instr(instr); + if (intr->intrinsic != nir_intrinsic_load_sample_mask_in) + return false; + + b->cursor = nir_before_instr(&intr->instr); nir_def *smp_msk = nir_ubitfield_extract_imm(b, nir_load_fs_meta_pco(b), 9, 16); - smp_msk = nir_iand(b, smp_msk, nir_load_savmsk_vm_pco(b)); - - return smp_msk; + nir_def_rewrite_uses(&intr->def, smp_msk); + nir_instr_remove(&intr->instr); + return true; } -static nir_def * -lower_color_write_enable(nir_builder *b, nir_instr *instr, UNUSED void *cb_data) +static bool lower_color_write_enable(nir_builder *b, + nir_intrinsic_instr *intr, + UNUSED void *cb_data) { - nir_intrinsic_instr *intr = nir_instr_as_intrinsic(instr); + if (intr->intrinsic != nir_intrinsic_store_output) + return false; + + gl_frag_result location = nir_intrinsic_io_semantics(intr).location; + if (!(location >= FRAG_RESULT_DATA0 && location < FRAG_RESULT_MAX)) + return false; nir_src *input_src = &intr->src[0]; nir_def *input = input_src->ssa; @@ -819,7 +808,7 @@ lower_color_write_enable(nir_builder *b, nir_instr *instr, UNUSED void *cb_data) nir_src_rewrite(input_src, nir_bcsel(b, color_write_enabled, input, prev_input)); - return NIR_LOWER_INSTR_PROGRESS; + return true; } /** @@ -853,28 +842,32 @@ bool pco_nir_pfo(nir_shader *shader, pco_fs_data *fs) * the stores */ if (!shader->info.internal) { - progress |= nir_shader_lower_instructions(shader, - is_frag_color_out, - lower_alpha_to_one, - &state); + progress |= nir_shader_intrinsics_pass(shader, + lower_alpha_to_one, + nir_metadata_control_flow, + &state); } - if (fs->meta_present.color_write_enable) - progress |= nir_shader_lower_instructions(shader, - is_frag_color_out, - lower_color_write_enable, - NULL); + if (fs->meta_present.color_write_enable) { + progress |= nir_shader_intrinsics_pass(shader, + lower_color_write_enable, + nir_metadata_control_flow, + NULL); + } - progress |= nir_shader_lower_instructions(shader, is_pfo, lower_pfo, &state); + progress |= nir_shader_intrinsics_pass(shader, + lower_pfo, + nir_metadata_control_flow, + &state); progress |= lower_isp_fb(&b, &state); progress |= sink_outputs(shader, &state); progress |= z_replicate(shader, &state); - progress |= nir_shader_lower_instructions(shader, - is_load_sample_mask, - lower_load_sample_mask, - NULL); + progress |= nir_shader_intrinsics_pass(shader, + lower_load_sample_mask, + nir_metadata_control_flow, + NULL); util_dynarray_fini(&state.stores); util_dynarray_fini(&state.loads); @@ -882,42 +875,33 @@ bool pco_nir_pfo(nir_shader *shader, pco_fs_data *fs) return progress; } -static nir_def *lower_pvi(nir_builder *b, nir_instr *instr, void *cb_data) +static bool lower_pvi(nir_builder *b, nir_intrinsic_instr *intr, void *cb_data) { - nir_intrinsic_instr *intr = nir_instr_as_intrinsic(instr); struct pvi_state *state = cb_data; - unsigned start_comp = nir_intrinsic_component(intr); - unsigned num_comps = intr->def.num_components; - - ASSERTED nir_src *offset = &intr->src[0]; - assert(nir_src_as_uint(*offset) == 0); - - struct nir_io_semantics io_semantics = nir_intrinsic_io_semantics(intr); - gl_vert_attrib location = io_semantics.location; - nir_def *attrib = state->attribs[location - VERT_ATTRIB_GENERIC0]; - assert(attrib); - - b->cursor = nir_before_instr(&intr->instr); - return nir_channels(b, attrib, BITFIELD_RANGE(start_comp, num_comps)); -} - -static bool is_pvi(const nir_instr *instr, const void *cb_data) -{ - if (instr->type != nir_instr_type_intrinsic) - return false; - - nir_intrinsic_instr *intr = nir_instr_as_intrinsic(instr); if (intr->intrinsic != nir_intrinsic_load_input) return false; if (is_processed(intr)) return false; + unsigned start_comp = nir_intrinsic_component(intr); + unsigned num_comps = intr->def.num_components; + ASSERTED gl_vert_attrib location = nir_intrinsic_io_semantics(intr).location; assert(location >= VERT_ATTRIB_GENERIC0 && location <= VERT_ATTRIB_GENERIC15); + ASSERTED nir_src *offset = &intr->src[0]; + assert(nir_src_as_uint(*offset) == 0); + + nir_def *attrib = state->attribs[location - VERT_ATTRIB_GENERIC0]; + assert(attrib); + + b->cursor = nir_before_instr(&intr->instr); + attrib = nir_channels(b, attrib, BITFIELD_RANGE(start_comp, num_comps)); + nir_def_rewrite_uses(&intr->def, attrib); + nir_instr_remove(&intr->instr); return true; } @@ -999,7 +983,10 @@ bool pco_nir_pvi(nir_shader *shader, pco_vs_data *vs) unpack_from_format(&b, packed_comps, base_type, format, 4); } - nir_shader_lower_instructions(shader, is_pvi, lower_pvi, &state); + nir_shader_intrinsics_pass(shader, + lower_pvi, + nir_metadata_control_flow, + &state); return true; } @@ -1074,24 +1061,7 @@ bool pco_nir_point_size(nir_shader *shader) return true; } -static bool is_fs_intr(const nir_instr *instr, UNUSED const void *cb_data) -{ - if (instr->type != nir_instr_type_intrinsic) - return false; - - nir_intrinsic_instr *intr = nir_instr_as_intrinsic(instr); - switch (intr->intrinsic) { - case nir_intrinsic_load_front_face: - return true; - - default: - break; - } - - return false; -} - -static nir_def *lower_front_face(nir_builder *b, nir_intrinsic_instr *intr) +static bool lower_front_face(nir_builder *b, nir_intrinsic_instr *intr) { nir_def *face_ccw = nir_load_face_ccw_pco(b); nir_def *front_face = nir_ieq_imm(b, face_ccw, 0); @@ -1103,32 +1073,35 @@ static nir_def *lower_front_face(nir_builder *b, nir_intrinsic_instr *intr) [PCO_FRONT_FACE_OP_TRUE] = nir_imm_true(b), }; - return nir_select_from_ssa_def_array(b, - ff_elems, - ARRAY_SIZE(ff_elems), - ff_op); + nir_def *ff_selected = + nir_select_from_ssa_def_array(b, ff_elems, ARRAY_SIZE(ff_elems), ff_op); + nir_def_rewrite_uses(&intr->def, ff_selected); + nir_instr_remove(&intr->instr); + return true; } -static nir_def * -lower_fs_intr(nir_builder *b, nir_instr *instr, UNUSED void *cb_data) +static bool +lower_fs_intr(nir_builder *b, nir_intrinsic_instr *intr, UNUSED void *cb_data) { - nir_intrinsic_instr *intr = nir_instr_as_intrinsic(instr); switch (intr->intrinsic) { case nir_intrinsic_load_front_face: + b->cursor = nir_before_instr(&intr->instr); return lower_front_face(b, intr); - default: break; } - UNREACHABLE(""); + return false; } bool pco_nir_lower_fs_intrinsics(nir_shader *shader) { assert(shader->info.stage == MESA_SHADER_FRAGMENT); - return nir_shader_lower_instructions(shader, is_fs_intr, lower_fs_intr, NULL); + return nir_shader_intrinsics_pass(shader, + lower_fs_intr, + nir_metadata_control_flow, + NULL); } static bool @@ -1232,13 +1205,9 @@ clone_clip_cull_stores(nir_builder *b, nir_intrinsic_instr *intr, void *data) return true; } -static bool is_clip_cull_load(const nir_instr *instr, - UNUSED const void *cb_data) +static bool +swap_clip_cull_load(nir_builder *b, nir_intrinsic_instr *intr, void *cb_data) { - if (instr->type != nir_instr_type_intrinsic) - return false; - - nir_intrinsic_instr *intr = nir_instr_as_intrinsic(instr); if (intr->intrinsic != nir_intrinsic_load_deref) return false; @@ -1248,16 +1217,11 @@ static bool is_clip_cull_load(const nir_instr *instr, nir_variable *var = nir_deref_instr_get_variable(deref); - return var->data.location == VARYING_SLOT_CLIP_DIST0 || - var->data.location == VARYING_SLOT_CLIP_DIST1; -} + if (var->data.location != VARYING_SLOT_CLIP_DIST0 && + var->data.location != VARYING_SLOT_CLIP_DIST1) + return false; -static nir_def * -swap_clip_cull_load(nir_builder *b, nir_instr *instr, void *cb_data) -{ - nir_intrinsic_instr *intr = nir_instr_as_intrinsic(instr); - nir_deref_instr *deref = nir_src_as_deref(intr->src[0]); - nir_variable *var = nir_deref_instr_get_variable(deref); + b->cursor = nir_before_instr(&intr->instr); unsigned var_index = var->data.location - VARYING_SLOT_CLIP_DIST0; nir_def *index = @@ -1265,7 +1229,10 @@ swap_clip_cull_load(nir_builder *b, nir_instr *instr, void *cb_data) index = nir_iadd_imm(b, index, var_index * 4); nir_variable *clone_var = cb_data; - return nir_load_array_var(b, clone_var, index); + nir_def *loaded_clone_var = nir_load_array_var(b, clone_var, index); + nir_def_rewrite_uses(&intr->def, loaded_clone_var); + nir_instr_remove(&intr->instr); + return true; } bool pco_nir_link_clip_cull_vars(nir_shader *producer, nir_shader *consumer) @@ -1304,10 +1271,10 @@ bool pco_nir_link_clip_cull_vars(nir_shader *producer, nir_shader *consumer) nir_variable_create(consumer, nir_var_shader_in, clone_var_type, NULL); clone_var->data.location = clone_slot; - nir_shader_lower_instructions(consumer, - is_clip_cull_load, - swap_clip_cull_load, - clone_var); + nir_shader_intrinsics_pass(consumer, + swap_clip_cull_load, + nir_metadata_control_flow, + clone_var); return true; } diff --git a/src/imagination/pco/pco_nir_sync.c b/src/imagination/pco/pco_nir_sync.c index 1aa3ea887b0..d915a794ebb 100644 --- a/src/imagination/pco/pco_nir_sync.c +++ b/src/imagination/pco/pco_nir_sync.c @@ -28,24 +28,29 @@ * \brief Lowers a barrier instruction. * * \param[in] b NIR builder. - * \param[in] instr NIR instruction. + * \param[in] intr NIR intrinsic instruction. * \param[in] cb_data User callback data. - * \return The replacement/lowered def. + * \return True if progress was made. */ -static nir_def *lower_barrier(nir_builder *b, nir_instr *instr, void *cb_data) +static bool +lower_barrier(nir_builder *b, nir_intrinsic_instr *intr, void *cb_data) { + if (intr->intrinsic != nir_intrinsic_barrier) + return false; + struct shader_info *info = &b->shader->info; bool *uses_usclib = cb_data; - nir_intrinsic_instr *intr = nir_instr_as_intrinsic(instr); mesa_scope exec_scope = nir_intrinsic_execution_scope(intr); unsigned wg_size = info->workgroup_size[0] * info->workgroup_size[1] * info->workgroup_size[2]; if (wg_size <= ROGUE_MAX_INSTANCES_PER_TASK || exec_scope == SCOPE_NONE || - exec_scope == SCOPE_SUBGROUP) - return NIR_LOWER_INSTR_PROGRESS_REPLACE; + exec_scope == SCOPE_SUBGROUP) { + nir_instr_remove(&intr->instr); + return true; + } /* TODO: We might be able to re-use barrier counters. */ unsigned counter_offset = info->shared_size; @@ -56,25 +61,10 @@ static nir_def *lower_barrier(nir_builder *b, nir_instr *instr, void *cb_data) unsigned num_slots = DIV_ROUND_UP(wg_size, ROGUE_MAX_INSTANCES_PER_TASK); - b->cursor = nir_before_instr(instr); + b->cursor = nir_before_instr(&intr->instr); usclib_barrier(b, nir_imm_int(b, num_slots), nir_imm_int(b, counter_offset)); - - return NIR_LOWER_INSTR_PROGRESS_REPLACE; -} - -/** - * \brief Filters barrier instructions. - * - * \param[in] instr NIR instruction. - * \param[in] cb_data User callback data. - * \return True if the instruction matches the filter. - */ -static bool is_barrier(const nir_instr *instr, UNUSED const void *cb_data) -{ - if (instr->type != nir_instr_type_intrinsic) - return false; - - return nir_instr_as_intrinsic(instr)->intrinsic == nir_intrinsic_barrier; + nir_instr_remove(&intr->instr); + return true; } /** @@ -85,23 +75,26 @@ static bool is_barrier(const nir_instr *instr, UNUSED const void *cb_data) */ bool pco_nir_lower_barriers(nir_shader *shader, pco_data *data) { - bool progress = nir_shader_lower_instructions(shader, - is_barrier, - lower_barrier, - &data->common.uses.usclib); + bool progress = nir_shader_intrinsics_pass(shader, + lower_barrier, + nir_metadata_none, + &data->common.uses.usclib); data->common.uses.barriers |= progress; return progress; } -static nir_def * -lower_usclib_atomic(nir_builder *b, nir_instr *instr, void *cb_data) +static bool +lower_usclib_atomic(nir_builder *b, nir_intrinsic_instr *intr, void *cb_data) { - nir_intrinsic_instr *intr = nir_instr_as_intrinsic(instr); + if (intr->intrinsic != nir_intrinsic_ssbo_atomic_swap && + intr->intrinsic != nir_intrinsic_global_atomic_swap_pco) + return false; + bool *uses_usclib = cb_data; - b->cursor = nir_before_instr(instr); + b->cursor = nir_before_instr(&intr->instr); if (intr->intrinsic == nir_intrinsic_ssbo_atomic_swap) { nir_def *buffer = intr->src[0].ssa; @@ -116,11 +109,11 @@ lower_usclib_atomic(nir_builder *b, nir_instr *instr, void *cb_data) assert(num_components == 1 && bit_size == 32); *uses_usclib = true; - return usclib_emu_ssbo_atomic_comp_swap(b, - buffer, - offset, - value, - value_swap); + nir_def *emulated = + usclib_emu_ssbo_atomic_comp_swap(b, buffer, offset, value, value_swap); + nir_def_rewrite_uses(&intr->def, emulated); + nir_instr_remove(&intr->instr); + return true; } nir_def *addr_data = intr->src[0].ssa; @@ -133,7 +126,12 @@ lower_usclib_atomic(nir_builder *b, nir_instr *instr, void *cb_data) assert(num_components == 1 && bit_size == 32); *uses_usclib = true; - return usclib_emu_global_atomic_comp_swap(b, addr, value, value_swap); + + nir_def *emulated = + usclib_emu_global_atomic_comp_swap(b, addr, value, value_swap); + nir_def_rewrite_uses(&intr->def, emulated); + nir_instr_remove(&intr->instr); + return true; } static bool lower_global_atomic_intrinsic(nir_builder *b, @@ -184,25 +182,6 @@ static bool lower_global_atomic_intrinsic(nir_builder *b, return true; } -/** - * \brief Filters atomic instructions emulated with usclib. - * - * \param[in] instr NIR instruction. - * \param[in] cb_data User callback data. - * \return True if the instruction matches the filter. - */ -static bool atomic_uses_usclib(const nir_instr *instr, - UNUSED const void *cb_data) -{ - if (instr->type != nir_instr_type_intrinsic) - return false; - - nir_intrinsic_instr *intr = nir_instr_as_intrinsic(instr); - - return intr->intrinsic == nir_intrinsic_ssbo_atomic_swap || - intr->intrinsic == nir_intrinsic_global_atomic_swap_pco; -} - /** * \brief Atomics lowering pass. * @@ -217,70 +196,58 @@ bool pco_nir_lower_atomics(nir_shader *shader, pco_data *data) lower_global_atomic_intrinsic, nir_metadata_none, NULL); - progress |= nir_shader_lower_instructions(shader, - atomic_uses_usclib, - lower_usclib_atomic, - &data->common.uses.usclib); + progress |= nir_shader_intrinsics_pass(shader, + lower_usclib_atomic, + nir_metadata_none, + &data->common.uses.usclib); return progress; } -static nir_def * -lower_subgroup_intrinsic(nir_builder *b, nir_instr *instr, void *cb_data) +static bool lower_subgroup_intrinsic(nir_builder *b, + nir_intrinsic_instr *intr, + void *cb_data) { - nir_intrinsic_instr *intr = nir_instr_as_intrinsic(instr); - assert(intr->def.num_components == 1); + nir_def *new_def; + + b->cursor = nir_before_instr(&intr->instr); switch (intr->intrinsic) { case nir_intrinsic_load_subgroup_size: - return nir_imm_int(b, 1); + new_def = nir_imm_int(b, 1); + break; case nir_intrinsic_load_subgroup_invocation: - return nir_imm_int(b, 0); + new_def = nir_imm_int(b, 0); + break; case nir_intrinsic_load_num_subgroups: - return nir_imm_int(b, - b->shader->info.workgroup_size[0] * - b->shader->info.workgroup_size[1] * - b->shader->info.workgroup_size[2]); + new_def = nir_imm_int(b, + b->shader->info.workgroup_size[0] * + b->shader->info.workgroup_size[1] * + b->shader->info.workgroup_size[2]); + break; case nir_intrinsic_load_subgroup_id: - return nir_load_local_invocation_index(b); + new_def = nir_load_local_invocation_index(b); + break; case nir_intrinsic_first_invocation: - return nir_imm_int(b, 0); + new_def = nir_imm_int(b, 0); + break; case nir_intrinsic_elect: - return nir_imm_true(b); + new_def = nir_imm_true(b); + break; default: - break; - } - - UNREACHABLE(""); -} - -static bool is_subgroup_intrinsic(const nir_instr *instr, - UNUSED const void *cb_data) -{ - if (instr->type != nir_instr_type_intrinsic) return false; - - nir_intrinsic_instr *intr = nir_instr_as_intrinsic(instr); - switch (intr->intrinsic) { - case nir_intrinsic_load_subgroup_size: - case nir_intrinsic_load_subgroup_invocation: - case nir_intrinsic_load_num_subgroups: - case nir_intrinsic_load_subgroup_id: - case nir_intrinsic_first_invocation: - case nir_intrinsic_elect: - return true; - - default: - break; } - return false; + nir_def_rewrite_uses(&intr->def, new_def); + nir_instr_remove(&intr->instr); + assert(intr->def.num_components == 1); + return true; } bool pco_nir_lower_subgroups(nir_shader *shader) @@ -289,8 +256,8 @@ bool pco_nir_lower_subgroups(nir_shader *shader) shader->info.min_subgroup_size = 1; shader->info.max_subgroup_size = 1; - return nir_shader_lower_instructions(shader, - is_subgroup_intrinsic, - lower_subgroup_intrinsic, - NULL); + return nir_shader_intrinsics_pass(shader, + lower_subgroup_intrinsic, + nir_metadata_control_flow, + NULL); } diff --git a/src/imagination/pco/pco_nir_tex.c b/src/imagination/pco/pco_nir_tex.c index b24bd0c90f8..fd07449c943 100644 --- a/src/imagination/pco/pco_nir_tex.c +++ b/src/imagination/pco/pco_nir_tex.c @@ -44,29 +44,36 @@ static inline nir_def *get_src_def(nir_tex_instr *tex, * \param[in] b NIR builder. * \param[in] tex NIR texture instruction. * \param[in] tex_state Texture state words. - * \return The replacement/lowered def. + * \return True if progress was made. */ -static nir_def *lower_tex_query_basic(nir_builder *b, - nir_tex_instr *tex, - nir_def *tex_state, - nir_def *tex_meta, - pco_data *data) +static bool lower_tex_query_basic(nir_builder *b, + nir_tex_instr *tex, + nir_def *tex_state, + nir_def *tex_meta, + pco_data *data) { + nir_def *new_def; + + b->cursor = nir_before_instr(&tex->instr); + switch (tex->op) { case nir_texop_query_levels: data->common.uses.usclib = true; - return usclib_tex_state_levels(b, tex_state); + new_def = usclib_tex_state_levels(b, tex_state); + break; case nir_texop_texture_samples: data->common.uses.usclib = true; - return usclib_tex_state_samples(b, tex_state); + new_def = usclib_tex_state_samples(b, tex_state); + break; case nir_texop_txs: { if (tex->sampler_dim == GLSL_SAMPLER_DIM_BUF) { assert(tex->def.num_components == 1); assert(!tex->is_array); - return nir_channel(b, tex_meta, PCO_IMAGE_META_BUFFER_ELEMS); + new_def = nir_channel(b, tex_meta, PCO_IMAGE_META_BUFFER_ELEMS); + break; } nir_def *num_comps = nir_imm_int(b, tex->def.num_components); @@ -84,14 +91,16 @@ static nir_def *lower_tex_query_basic(nir_builder *b, data->common.uses.usclib = true; - return nir_trim_vector(b, size_comps, tex->def.num_components); - } - - default: + new_def = nir_trim_vector(b, size_comps, tex->def.num_components); break; } - UNREACHABLE(""); + default: + UNREACHABLE(""); + } + + nir_def_rewrite_uses(&tex->def, new_def); + return true; } static inline enum pco_dim to_pco_dim(enum glsl_sampler_dim dim) @@ -442,13 +451,12 @@ static nir_def *lower_tex_shadow(nir_builder *b, * \brief Lowers a texture instruction. * * \param[in] b NIR builder. - * \param[in] instr NIR instruction. + * \param[in] instr NIR texture instruction. * \param[in] cb_data User callback data. - * \return The replacement/lowered def. + * \return True if progress was made. */ -static nir_def *lower_tex(nir_builder *b, nir_instr *instr, void *cb_data) +static bool lower_tex(nir_builder *b, nir_tex_instr *tex, void *cb_data) { - nir_tex_instr *tex = nir_instr_as_tex(instr); struct state *state = cb_data; pco_data *data = state->data; pco_ctx *ctx = state->ctx; @@ -468,7 +476,7 @@ static nir_def *lower_tex(nir_builder *b, nir_instr *instr, void *cb_data) ? PVR_HAS_FEATURE(dev_info, tpu_extended_integer_lookup) : false; - b->cursor = nir_before_instr(instr); + b->cursor = nir_before_instr(&tex->instr); /* Process tex sources, build up the smp flags and data words. */ BITSET_DECLARE(tex_src_set, nir_num_tex_src_types) = { 0 }; @@ -771,20 +779,8 @@ static nir_def *lower_tex(nir_builder *b, nir_instr *instr, void *cb_data) result = lower_tex_shadow(b, result, comparator, compare_op); } - - return result; -} - -/** - * \brief Filters texture instructions. - * - * \param[in] instr NIR instruction. - * \param[in] cb_data User callback data. - * \return True if the instruction matches the filter. - */ -static bool is_tex(const nir_instr *instr, UNUSED const void *cb_data) -{ - return instr->type == nir_instr_type_tex; + nir_def_rewrite_uses(&tex->def, result); + return true; } /** @@ -802,7 +798,10 @@ bool pco_nir_lower_tex(nir_shader *shader, pco_data *data, pco_ctx *ctx) .ctx = ctx, }; - return nir_shader_lower_instructions(shader, is_tex, lower_tex, &state); + return nir_shader_tex_pass(shader, + lower_tex, + nir_metadata_control_flow, + &state); } static enum util_format_type nir_type_to_util_type(nir_alu_type nir_type) @@ -838,19 +837,15 @@ static enum pipe_format nir_type_to_pipe_format(nir_alu_type nir_type, pure_integer); } -static nir_def *lower_image(nir_builder *b, nir_instr *instr, void *cb_data) +static bool +lower_image(nir_builder *b, nir_intrinsic_instr *intr, void *cb_data) { - nir_intrinsic_instr *intr = nir_instr_as_intrinsic(instr); struct state *state = cb_data; pco_data *data = state->data; pco_ctx *ctx = state->ctx; const struct pvr_device_info *dev_info = ctx->dev_info; - enum glsl_sampler_dim image_dim = nir_intrinsic_image_dim(intr); - bool is_array = nir_intrinsic_image_array(intr); - enum pipe_format format = nir_intrinsic_format(intr); - unsigned desc_set = nir_src_comp_as_uint(intr->src[0], 0); - unsigned binding = nir_src_comp_as_uint(intr->src[0], 1); - nir_def *elem = nir_channel(b, intr->src[0].ssa, 2); + + b->cursor = nir_before_instr(&intr->instr); nir_def *lod = NULL; switch (intr->intrinsic) { @@ -872,9 +867,16 @@ static nir_def *lower_image(nir_builder *b, nir_instr *instr, void *cb_data) break; default: - UNREACHABLE(""); + return false; } + enum glsl_sampler_dim image_dim = nir_intrinsic_image_dim(intr); + bool is_array = nir_intrinsic_image_array(intr); + enum pipe_format format = nir_intrinsic_format(intr); + unsigned desc_set = nir_src_comp_as_uint(intr->src[0], 0); + unsigned binding = nir_src_comp_as_uint(intr->src[0], 1); + nir_def *elem = nir_channel(b, intr->src[0].ssa, 2); + if (intr->intrinsic == nir_intrinsic_image_deref_size) { if (image_dim == GLSL_SAMPLER_DIM_BUF) { assert(intr->def.num_components == 1); @@ -883,8 +885,11 @@ static nir_def *lower_image(nir_builder *b, nir_instr *instr, void *cb_data) elem, .desc_set = desc_set, .binding = binding); - - return nir_channel(b, tex_meta, PCO_IMAGE_META_BUFFER_ELEMS); + nir_def *buf_elems = + nir_channel(b, tex_meta, PCO_IMAGE_META_BUFFER_ELEMS); + nir_def_rewrite_uses(&intr->def, buf_elems); + nir_instr_remove(&intr->instr); + return true; } nir_def *tex_state = nir_load_tex_state_pco(b, @@ -907,7 +912,11 @@ static nir_def *lower_image(nir_builder *b, nir_instr *instr, void *cb_data) data->common.uses.usclib = true; - return nir_trim_vector(b, size_comps, intr->def.num_components); + nir_def *image_size = + nir_trim_vector(b, size_comps, intr->def.num_components); + nir_def_rewrite_uses(&intr->def, image_size); + nir_instr_remove(&intr->instr); + return true; } nir_alu_type type = nir_type_invalid; @@ -1071,22 +1080,24 @@ static nir_def *lower_image(nir_builder *b, nir_instr *instr, void *cb_data) if (ia) { assert(!is_array); - nir_load_const_instr *load = - nir_def_as_load_const(intr->src[0].ssa); + nir_load_const_instr *load = nir_def_as_load_const(intr->src[0].ssa); bool onchip = load->def.num_components == 4; if (onchip) { unsigned ia_idx = nir_src_comp_as_uint(intr->src[0], 3); - return nir_load_output(b, - intr->def.num_components, - intr->def.bit_size, - nir_imm_int(b, 0), - .base = ia_idx, - .component = 0, - .dest_type = nir_intrinsic_dest_type(intr), - .io_semantics.location = FRAG_RESULT_COLOR, - .io_semantics.num_slots = 1/*, - .io_semantics.fb_fetch_output = true*/); + nir_def *loaded_ia = nir_load_output(b, + intr->def.num_components, + intr->def.bit_size, + nir_imm_int(b, 0), + .base = ia_idx, + .component = 0, + .dest_type = nir_intrinsic_dest_type(intr), + .io_semantics.location = FRAG_RESULT_COLOR, + .io_semantics.num_slots = 1/*, + .io_semantics.fb_fetch_output = true*/); + nir_def_rewrite_uses(&intr->def, loaded_ia); + nir_instr_remove(&intr->instr); + return true; } } @@ -1210,10 +1221,13 @@ static nir_def *lower_image(nir_builder *b, nir_instr *instr, void *cb_data) nir_def *addr_data = nir_vec4(b, addr_lo, addr_hi, compare, dma_data); - return nir_global_atomic_swap_pco(b, - addr_data, - .atomic_op = - nir_intrinsic_atomic_op(intr)); + nir_def *atomic_swap = nir_global_atomic_swap_pco( + b, + addr_data, + .atomic_op = nir_intrinsic_atomic_op(intr)); + nir_def_rewrite_uses(&intr->def, atomic_swap); + nir_instr_remove(&intr->instr); + return true; } nir_def *dma_data = intr->src[3].ssa; @@ -1222,9 +1236,13 @@ static nir_def *lower_image(nir_builder *b, nir_instr *instr, void *cb_data) data->common.uses.usclib = true; - return nir_global_atomic_pco(b, - addr_data, - .atomic_op = nir_intrinsic_atomic_op(intr)); + nir_def *atomic = + nir_global_atomic_pco(b, + addr_data, + .atomic_op = nir_intrinsic_atomic_op(intr)); + nir_def_rewrite_uses(&intr->def, atomic); + nir_instr_remove(&intr->instr); + return true; } unsigned smp_desc = ia ? PCO_IA_SAMPLER : PCO_POINT_SAMPLER; @@ -1383,31 +1401,14 @@ static nir_def *lower_image(nir_builder *b, nir_instr *instr, void *cb_data) nir_intrinsic_instr *smp = pco_emit_nir_smp(b, ¶ms); - if (intr->intrinsic == nir_intrinsic_image_deref_load) - return &smp->def; - - return NIR_LOWER_INSTR_PROGRESS_REPLACE; -} - -static bool is_image(const nir_instr *instr, UNUSED const void *cb_data) -{ - if (instr->type != nir_instr_type_intrinsic) - return false; - - nir_intrinsic_instr *intr = nir_instr_as_intrinsic(instr); - switch (intr->intrinsic) { - case nir_intrinsic_image_deref_load: - case nir_intrinsic_image_deref_store: - case nir_intrinsic_image_deref_atomic: - case nir_intrinsic_image_deref_atomic_swap: - case nir_intrinsic_image_deref_size: + if (intr->intrinsic == nir_intrinsic_image_deref_load) { + nir_def_rewrite_uses(&intr->def, &smp->def); + nir_instr_remove(&intr->instr); return true; - - default: - break; } - return false; + nir_instr_remove(&intr->instr); + return true; } bool pco_nir_lower_images(nir_shader *shader, pco_data *data, pco_ctx *ctx) @@ -1416,5 +1417,9 @@ bool pco_nir_lower_images(nir_shader *shader, pco_data *data, pco_ctx *ctx) .data = data, .ctx = ctx, }; - return nir_shader_lower_instructions(shader, is_image, lower_image, &state); + + return nir_shader_intrinsics_pass(shader, + lower_image, + nir_metadata_none, + &state); } diff --git a/src/imagination/pco/pco_nir_vk.c b/src/imagination/pco/pco_nir_vk.c index 1cd7910bb48..e7818bb4566 100644 --- a/src/imagination/pco/pco_nir_vk.c +++ b/src/imagination/pco/pco_nir_vk.c @@ -41,11 +41,11 @@ set_resource_used(pco_common_data *common, unsigned desc_set, unsigned binding) * \param[in] b NIR builder. * \param[in] intr NIR intrinsic instruction. * \param[in] common Shader common data. - * \return The replacement/lowered def. + * \return True if progress was made. */ -static nir_def *lower_load_vulkan_descriptor(nir_builder *b, - nir_intrinsic_instr *intr, - pco_common_data *common) +static bool lower_load_vulkan_descriptor(nir_builder *b, + nir_intrinsic_instr *intr, + pco_common_data *common) { nir_intrinsic_instr *vk_res_idx = nir_src_as_intrinsic(intr->src[0]); assert(vk_res_idx->intrinsic == nir_intrinsic_vulkan_resource_index); @@ -58,8 +58,13 @@ static nir_def *lower_load_vulkan_descriptor(nir_builder *b, set_resource_used(common, desc_set, binding); + b->cursor = nir_before_instr(&intr->instr); + uint32_t desc_set_binding = pco_pack_desc(desc_set, binding); - return nir_imm_ivec3(b, desc_set_binding, elem, 0); + nir_def *desc_ref = nir_imm_ivec3(b, desc_set_binding, elem, 0); + nir_def_rewrite_uses(&intr->def, desc_ref); + nir_instr_remove(&intr->instr); + return true; } static nir_def *array_elem_from_deref(nir_builder *b, nir_deref_instr *deref) @@ -88,14 +93,13 @@ static inline bool is_comb_img_smp(unsigned desc_set, return binding_data->is_img_smp; } -static void lower_tex_deref_to_binding(nir_builder *b, +static bool lower_tex_deref_to_binding(nir_builder *b, nir_tex_instr *tex, unsigned deref_index, pco_common_data *common) { nir_tex_src *deref_src = &tex->src[deref_index]; - nir_deref_instr *deref = - nir_def_as_deref(deref_src->src.ssa); + nir_deref_instr *deref = nir_def_as_deref(deref_src->src.ssa); b->cursor = nir_before_instr(&tex->instr); @@ -117,9 +121,10 @@ static void lower_tex_deref_to_binding(nir_builder *b, } nir_src_rewrite(&deref_src->src, elem); + return true; } -static void +static bool add_txf_sampler(nir_builder *b, nir_tex_instr *tex, pco_common_data *common) { int deref_index = nir_tex_instr_src_index(tex, nir_tex_src_backend1); @@ -142,25 +147,29 @@ add_txf_sampler(nir_builder *b, nir_tex_instr *tex, pco_common_data *common) tex->sampler_index = pco_pack_desc(desc_set, binding); nir_tex_instr_add_src(tex, nir_tex_src_backend2, elem); + return true; } -static inline void +static inline bool lower_tex_derefs(nir_builder *b, nir_tex_instr *tex, pco_common_data *common) { int deref_index; + bool progress = false; deref_index = nir_tex_instr_src_index(tex, nir_tex_src_texture_deref); if (deref_index >= 0) - lower_tex_deref_to_binding(b, tex, deref_index, common); + progress |= lower_tex_deref_to_binding(b, tex, deref_index, common); deref_index = nir_tex_instr_src_index(tex, nir_tex_src_sampler_deref); if (deref_index >= 0) - lower_tex_deref_to_binding(b, tex, deref_index, common); + progress |= lower_tex_deref_to_binding(b, tex, deref_index, common); else if (tex->op == nir_texop_txf || tex->op == nir_texop_txf_ms) - add_txf_sampler(b, tex, common); + progress |= add_txf_sampler(b, tex, common); + + return progress; } -static nir_def * +static bool lower_image_derefs(nir_builder *b, nir_intrinsic_instr *intr, pco_data *data) { nir_src *deref_src = &intr->src[0]; @@ -200,7 +209,7 @@ lower_image_derefs(nir_builder *b, nir_intrinsic_instr *intr, pco_data *data) nir_src_rewrite(deref_src, index); - return NIR_LOWER_INSTR_PROGRESS; + return true; } /* Sampler not needed for on-chip input attachments. */ @@ -217,18 +226,17 @@ lower_image_derefs(nir_builder *b, nir_intrinsic_instr *intr, pco_data *data) nir_src_rewrite(deref_src, index); - return NIR_LOWER_INSTR_PROGRESS; + return true; } -static nir_def *lower_is_null_descriptor(nir_builder *b, - nir_intrinsic_instr *intr) +static bool lower_is_null_descriptor(nir_builder *b, nir_intrinsic_instr *intr) { nir_src *deref_src = &intr->src[0]; nir_deref_instr *deref = nir_src_as_deref(*deref_src); /* Will be taken care of by lower_load_vulkan_descriptor. */ if (!deref) - return NULL; + return false; b->cursor = nir_before_instr(&intr->instr); @@ -242,102 +250,61 @@ static nir_def *lower_is_null_descriptor(nir_builder *b, nir_def *index = nir_vec2(b, nir_imm_int(b, desc_set_binding), elem); nir_src_rewrite(deref_src, index); - return NIR_LOWER_INSTR_PROGRESS; + return true; } /** - * \brief Lowers a Vulkan-related instruction. + * \brief Lowers a Vulkan-related texture instruction. * * \param[in] b NIR builder. - * \param[in] instr NIR instruction. + * \param[in] tex NIR texture instruction. * \param[in] cb_data User callback data. - * \return The replacement/lowered def. + * \return True if progress was made. */ -static nir_def *lower_vk(nir_builder *b, nir_instr *instr, void *cb_data) +static bool lower_vk_tex(nir_builder *b, nir_tex_instr *tex, void *cb_data) { pco_data *data = cb_data; pco_common_data *common = &data->common; - - switch (instr->type) { - case nir_instr_type_intrinsic: { - nir_intrinsic_instr *intr = nir_instr_as_intrinsic(instr); - switch (intr->intrinsic) { - case nir_intrinsic_load_vulkan_descriptor: - return lower_load_vulkan_descriptor(b, intr, common); - - case nir_intrinsic_image_deref_load: - case nir_intrinsic_image_deref_store: - case nir_intrinsic_image_deref_atomic: - case nir_intrinsic_image_deref_atomic_swap: - case nir_intrinsic_image_deref_size: - return lower_image_derefs(b, intr, data); - - case nir_intrinsic_is_null_descriptor: - return lower_is_null_descriptor(b, intr); - - default: - break; - } - - break; + if (nir_tex_instr_src_index(tex, nir_tex_src_texture_deref) >= 0 || + nir_tex_instr_src_index(tex, nir_tex_src_sampler_deref) >= 0) { + return lower_tex_derefs(b, tex, common); } - - case nir_instr_type_tex: { - nir_tex_instr *tex = nir_instr_as_tex(instr); - lower_tex_derefs(b, tex, common); - return NIR_LOWER_INSTR_PROGRESS; - } - - default: - break; - } - - UNREACHABLE(""); + return false; } /** - * \brief Filters Vulkan-related instructions. + * \brief Lowers a Vulkan-related intrinsic instruction. * - * \param[in] instr NIR instruction. + * \param[in] b NIR builder. + * \param[in] intr NIR intrinsic instruction. * \param[in] cb_data User callback data. - * \return True if the instruction matches the filter. + * \return True if progress was made. */ -static bool is_vk(const nir_instr *instr, UNUSED const void *cb_data) +static bool +lower_vk_intr(nir_builder *b, nir_intrinsic_instr *intr, void *cb_data) { - switch (instr->type) { - case nir_instr_type_intrinsic: { - nir_intrinsic_instr *intr = nir_instr_as_intrinsic(instr); - switch (intr->intrinsic) { - case nir_intrinsic_load_vulkan_descriptor: - case nir_intrinsic_is_null_descriptor: - case nir_intrinsic_image_deref_load: - case nir_intrinsic_image_deref_store: - case nir_intrinsic_image_deref_atomic: - case nir_intrinsic_image_deref_atomic_swap: - case nir_intrinsic_image_deref_size: - return true; + pco_data *data = cb_data; + pco_common_data *common = &data->common; + + b->cursor = nir_before_instr(&intr->instr); - default: - break; - } + switch (intr->intrinsic) { + case nir_intrinsic_load_vulkan_descriptor: + return lower_load_vulkan_descriptor(b, intr, common); - break; - } + case nir_intrinsic_image_deref_load: + case nir_intrinsic_image_deref_store: + case nir_intrinsic_image_deref_atomic: + case nir_intrinsic_image_deref_atomic_swap: + case nir_intrinsic_image_deref_size: + return lower_image_derefs(b, intr, data); - case nir_instr_type_tex: { - nir_tex_instr *tex = nir_instr_as_tex(instr); - if (nir_tex_instr_src_index(tex, nir_tex_src_texture_deref) >= 0 || - nir_tex_instr_src_index(tex, nir_tex_src_sampler_deref) >= 0) { - return true; - } - - FALLTHROUGH; - } + case nir_intrinsic_is_null_descriptor: + return lower_is_null_descriptor(b, intr); default: break; } - return false; } @@ -352,7 +319,14 @@ bool pco_nir_lower_vk(nir_shader *shader, pco_data *data) { bool progress = false; - progress |= nir_shader_lower_instructions(shader, is_vk, lower_vk, data); + progress |= nir_shader_intrinsics_pass(shader, + lower_vk_intr, + nir_metadata_control_flow, + data); + progress |= nir_shader_tex_pass(shader, + lower_vk_tex, + nir_metadata_control_flow, + data); return progress; }