treewide: use nir_def_replace sometimes

Two Coccinelle patches here. Didn't catch nearly as much as I would've liked but
it's a start.

Coccinelle patch:

    @@
    expression intr, repl;
    @@

    -nir_def_rewrite_uses(&intr->def, repl);
    -nir_instr_remove(&intr->instr);
    +nir_def_replace(&intr->def, repl);

Coccinelle patch:

    @@
    identifier intr;
    expression instr, repl;
    @@

    nir_intrinsic_instr *intr = nir_instr_as_intrinsic(instr);
    ...
    -nir_def_rewrite_uses(&intr->def, repl);
    -nir_instr_remove(instr);
    +nir_def_replace(&intr->def, repl);

Signed-off-by: Alyssa Rosenzweig <alyssa@rosenzweig.io>
Reviewed-by: Faith Ekstrand <faith.ekstrand@collabora.com>
Reviewed-by: Juan A. Suarez Romero <jasuarez@igalia.com> [broadcom]
Reviewed-by: Vasily Khoruzhick <anarsoul@gmail.com> [lima]
Reviewed-by: Christian Gmeiner <cgmeiner@igalia.com> [etna]
Reviewed-by: Pavel Ondračka <pavel.ondracka@gmail.com> [r300]
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/29817>
This commit is contained in:
Alyssa Rosenzweig 2024-06-20 12:07:26 -04:00 committed by Marge Bot
parent bbdd34b4ad
commit da752ed7c1
101 changed files with 180 additions and 416 deletions

View file

@ -174,8 +174,7 @@ lower_intrinsic_to_arg(nir_builder *b, nir_instr *instr, void *state)
}
assert(replacement);
nir_def_rewrite_uses(&intrin->def, replacement);
nir_instr_remove(&intrin->instr);
nir_def_replace(&intrin->def, replacement);
return true;
}
@ -1475,8 +1474,7 @@ split_pack_half(nir_builder *b, nir_instr *instr, void *param)
*/
nir_def *lo = nir_f2f16(b, nir_ssa_for_alu_src(b, alu, 0));
nir_def *hi = nir_f2f16(b, nir_ssa_for_alu_src(b, alu, 1));
nir_def_rewrite_uses(&alu->def, nir_pack_32_2x16_split(b, lo, hi));
nir_instr_remove(&alu->instr);
nir_def_replace(&alu->def, nir_pack_32_2x16_split(b, lo, hi));
return true;
}

View file

@ -187,9 +187,7 @@ lower_ps_load_barycentric(nir_builder *b, nir_intrinsic_instr *intrin, lower_ps_
b->cursor = nir_before_instr(&intrin->instr);
nir_def *replacement = nir_load_var(b, var);
nir_def_rewrite_uses(&intrin->def, replacement);
nir_instr_remove(&intrin->instr);
nir_def_replace(&intrin->def, replacement);
return true;
}
@ -261,9 +259,7 @@ lower_ps_load_sample_mask_in(nir_builder *b, nir_intrinsic_instr *intrin, lower_
nir_def *sample_mask = nir_load_sample_mask_in(b);
nir_def *replacement = nir_iand(b, sample_mask, submask);
nir_def_rewrite_uses(&intrin->def, replacement);
nir_instr_remove(&intrin->instr);
nir_def_replace(&intrin->def, replacement);
return true;
}

View file

@ -188,8 +188,7 @@ visit_get_ssbo_size(nir_builder *b, apply_layout_state *state, nir_intrinsic_ins
size = nir_channel(b, desc, 2);
}
nir_def_rewrite_uses(&intrin->def, size);
nir_instr_remove(&intrin->instr);
nir_def_replace(&intrin->def, size);
}
static nir_def *
@ -335,8 +334,7 @@ update_image_intrinsic(nir_builder *b, apply_layout_state *state, nir_intrinsic_
nir_intrinsic_access(intrin) & ACCESS_NON_UNIFORM, NULL, !is_load);
if (intrin->intrinsic == nir_intrinsic_image_deref_descriptor_amd) {
nir_def_rewrite_uses(&intrin->def, desc);
nir_instr_remove(&intrin->instr);
nir_def_replace(&intrin->def, desc);
} else {
nir_rewrite_image_intrinsic(intrin, desc, true);
}
@ -437,8 +435,7 @@ apply_layout_to_intrin(nir_builder *b, apply_layout_state *state, nir_intrinsic_
update_image_intrinsic(b, state, intrin);
break;
case nir_intrinsic_load_push_constant: {
nir_def_rewrite_uses(&intrin->def, load_push_constant(b, state, intrin));
nir_instr_remove(&intrin->instr);
nir_def_replace(&intrin->def, load_push_constant(b, state, intrin));
break;
}
default:
@ -513,8 +510,7 @@ apply_layout_to_tex(nir_builder *b, apply_layout_state *state, nir_tex_instr *te
}
if (tex->op == nir_texop_descriptor_amd) {
nir_def_rewrite_uses(&tex->def, image);
nir_instr_remove(&tex->instr);
nir_def_replace(&tex->def, image);
return;
}

View file

@ -234,8 +234,7 @@ lower_load_barycentric_coord(nir_builder *b, lower_fs_barycentric_state *state,
}
}
nir_def_rewrite_uses(&intrin->def, new_dest);
nir_instr_remove(&intrin->instr);
nir_def_replace(&intrin->def, new_dest);
return true;
}

View file

@ -176,8 +176,7 @@ radv_nir_lower_draw_id_to_zero_callback(struct nir_builder *b, nir_intrinsic_ins
return false;
nir_def *replacement = nir_imm_zero(b, intrin->def.num_components, intrin->def.bit_size);
nir_def_rewrite_uses(&intrin->def, replacement);
nir_instr_remove(&intrin->instr);
nir_def_replace(&intrin->def, replacement);
nir_instr_free(&intrin->instr);
return true;

View file

@ -385,8 +385,7 @@ lower_vs_input_instr(nir_builder *b, nir_intrinsic_instr *intrin, void *state)
replacement = lower_load_vs_input(b, intrin, s);
}
nir_def_rewrite_uses(&intrin->def, replacement);
nir_instr_remove(&intrin->instr);
nir_def_replace(&intrin->def, replacement);
nir_instr_free(&intrin->instr);
return true;

View file

@ -151,8 +151,7 @@ lower_rt_derefs(nir_shader *shader)
b.cursor = nir_before_instr(&deref->instr);
nir_deref_instr *replacement =
nir_build_deref_cast(&b, arg_offset, nir_var_function_temp, deref->var->type, 0);
nir_def_rewrite_uses(&deref->def, &replacement->def);
nir_instr_remove(&deref->instr);
nir_def_replace(&deref->def, &replacement->def);
}
}
}
@ -1091,13 +1090,11 @@ lower_any_hit_for_intersection(nir_shader *any_hit)
break;
case nir_intrinsic_load_ray_t_max:
nir_def_rewrite_uses(&intrin->def, hit_t);
nir_instr_remove(&intrin->instr);
nir_def_replace(&intrin->def, hit_t);
break;
case nir_intrinsic_load_ray_hit_kind:
nir_def_rewrite_uses(&intrin->def, hit_kind);
nir_instr_remove(&intrin->instr);
nir_def_replace(&intrin->def, hit_kind);
break;
/* We place all any_hit scratch variables after intersection scratch variables.

View file

@ -20,8 +20,7 @@ lower_to_per_sample(nir_builder *b, nir_intrinsic_instr *intr, void *data)
case nir_intrinsic_load_sample_id: {
nir_def *mask = nir_u2u32(b, nir_load_active_samples_agx(b));
nir_def *bit = nir_ufind_msb(b, mask);
nir_def_rewrite_uses(&intr->def, nir_u2uN(b, bit, intr->def.bit_size));
nir_instr_remove(&intr->instr);
nir_def_replace(&intr->def, nir_u2uN(b, bit, intr->def.bit_size));
return true;
}

View file

@ -30,8 +30,7 @@ lower(nir_builder *b, nir_intrinsic_instr *intr, void *data)
/* Handle the center special case */
if (!b->shader->info.fs.uses_sample_shading) {
assert(intr->intrinsic == nir_intrinsic_load_sample_pos_or_center);
nir_def_rewrite_uses(&intr->def, nir_imm_vec2(b, 0.5, 0.5));
nir_instr_remove(&intr->instr);
nir_def_replace(&intr->def, nir_imm_vec2(b, 0.5, 0.5));
return true;
}
@ -65,8 +64,7 @@ lower(nir_builder *b, nir_intrinsic_instr *intr, void *data)
}
/* Collect and rewrite */
nir_def_rewrite_uses(&intr->def, nir_vec2(b, xy[0], xy[1]));
nir_instr_remove(&intr->instr);
nir_def_replace(&intr->def, nir_vec2(b, xy[0], xy[1]));
return true;
}

View file

@ -240,8 +240,7 @@ lower_tes(nir_builder *b, nir_intrinsic_instr *intr, void *data)
nir_def *repl = lower_tes_impl(b, intr, data);
if (repl) {
nir_def_rewrite_uses(&intr->def, repl);
nir_instr_remove(&intr->instr);
nir_def_replace(&intr->def, repl);
return true;
} else {
return false;

View file

@ -203,8 +203,7 @@ lower_input_to_prolog(nir_builder *b, nir_intrinsic_instr *intr, void *data)
BITSET_SET(comps_read, base + c);
}
nir_def_rewrite_uses(&intr->def, val);
nir_instr_remove(&intr->instr);
nir_def_replace(&intr->def, val);
return true;
}

View file

@ -103,9 +103,7 @@ lower_load_bitsize(nir_builder *b,
}
nir_def *new_dst = nir_vec(b, dest_components, num_comp);
nir_def_rewrite_uses(&intr->def, new_dst);
nir_instr_remove(&intr->instr);
nir_def_replace(&intr->def, new_dst);
return true;
}

View file

@ -1536,8 +1536,7 @@ lower_load_num_subgroups(struct v3d_compile *c,
c->s->info.workgroup_size[1] *
c->s->info.workgroup_size[2], V3D_CHANNELS);
nir_def *result = nir_imm_int(b, num_subgroups);
nir_def_rewrite_uses(&intr->def, result);
nir_instr_remove(&intr->instr);
nir_def_replace(&intr->def, result);
}
static bool

View file

@ -573,9 +573,7 @@ lower_vulkan_resource_index(nir_builder *b,
* vulkan_load_descriptor return a vec2 providing an index and
* offset. Our backend compiler only cares about the index part.
*/
nir_def_rewrite_uses(&instr->def,
nir_imm_ivec2(b, index, 0));
nir_instr_remove(&instr->instr);
nir_def_replace(&instr->def, nir_imm_ivec2(b, index, 0));
}
static uint8_t
@ -830,8 +828,7 @@ lower_intrinsic(nir_builder *b,
/* Loading the descriptor happens as part of load/store instructions,
* so for us this is a no-op.
*/
nir_def_rewrite_uses(&instr->def, instr->src[0].ssa);
nir_instr_remove(&instr->instr);
nir_def_replace(&instr->def, instr->src[0].ssa);
return true;
}

View file

@ -3457,9 +3457,7 @@ replace_unused_interpolate_at_with_undef(nir_builder *b, nir_instr *instr,
nir_def *undef =
nir_undef(b, intrin->def.num_components,
intrin->def.bit_size);
nir_def_rewrite_uses(&intrin->def, undef);
nir_instr_remove(&intrin->instr);
nir_def_replace(&intrin->def, undef);
return true;
}
}

View file

@ -1163,9 +1163,7 @@ opt_remove_sampler_cast(nir_deref_instr *cast)
/* We're a cast from a more detailed sampler type to a bare sampler or a
* texture type with the same dimensionality.
*/
nir_def_rewrite_uses(&cast->def,
&parent->def);
nir_instr_remove(&cast->instr);
nir_def_replace(&cast->def, &parent->def);
/* Recursively crawl the deref tree and clean up types */
nir_deref_instr_fixup_child_types(parent);
@ -1284,9 +1282,7 @@ opt_deref_ptr_as_array(nir_builder *b, nir_deref_instr *deref)
parent->cast.align_mul == 0 &&
nir_deref_cast_is_trivial(parent))
parent = nir_deref_instr_parent(parent);
nir_def_rewrite_uses(&deref->def,
&parent->def);
nir_instr_remove(&deref->instr);
nir_def_replace(&deref->def, &parent->def);
return true;
}
@ -1464,8 +1460,7 @@ opt_known_deref_mode_is(nir_builder *b, nir_intrinsic_instr *intrin)
if (deref_is == NULL)
return false;
nir_def_rewrite_uses(&intrin->def, deref_is);
nir_instr_remove(&intrin->instr);
nir_def_replace(&intrin->def, deref_is);
return true;
}

View file

@ -104,13 +104,7 @@ nir_inline_function_impl(struct nir_builder *b,
unsigned param_idx = nir_intrinsic_param_idx(load);
assert(param_idx < impl->function->num_params);
nir_def_rewrite_uses(&load->def,
params[param_idx]);
/* Remove any left-over load_param intrinsics because they're soon
* to be in another function and therefore no longer valid.
*/
nir_instr_remove(&load->instr);
nir_def_replace(&load->def, params[param_idx]);
break;
}

View file

@ -414,8 +414,7 @@ nir_inline_uniforms(nir_shader *shader, unsigned num_uniforms,
if (offset == uniform_dw_offsets[i]) {
b.cursor = nir_before_instr(&intr->instr);
nir_def *def = nir_imm_int(&b, uniform_values[i]);
nir_def_rewrite_uses(&intr->def, def);
nir_instr_remove(&intr->instr);
nir_def_replace(&intr->def, def);
break;
}
}
@ -458,9 +457,8 @@ nir_inline_uniforms(nir_shader *shader, unsigned num_uniforms,
}
/* Replace the original uniform load. */
nir_def_rewrite_uses(&intr->def,
nir_vec(&b, components, num_components));
nir_instr_remove(&intr->instr);
nir_def_replace(&intr->def,
nir_vec(&b, components, num_components));
}
}
}

View file

@ -214,8 +214,7 @@ lower_alu_instr(nir_builder *b, nir_instr *instr_, UNUSED void *cb_data)
}
if (lowered) {
nir_def_rewrite_uses(&instr->def, lowered);
nir_instr_remove(&instr->instr);
nir_def_replace(&instr->def, lowered);
return true;
} else {
return false;

View file

@ -113,9 +113,7 @@ nir_lower_array_deref_of_vec_impl(nir_function_impl *impl,
nir_def *scalar =
nir_vector_extract(&b, &intrin->def, index);
if (scalar->parent_instr->type == nir_instr_type_undef) {
nir_def_rewrite_uses(&intrin->def,
scalar);
nir_instr_remove(&intrin->instr);
nir_def_replace(&intrin->def, scalar);
} else {
nir_def_rewrite_uses_after(&intrin->def,
scalar,

View file

@ -338,8 +338,7 @@ split_phi(nir_builder *b, nir_phi_instr *phi)
b->cursor = nir_after_phis(nir_cursor_current_block(b->cursor));
nir_def *merged = nir_pack_64_2x32_split(b, &lowered[0]->def, &lowered[1]->def);
nir_def_rewrite_uses(&phi->def, merged);
nir_instr_remove(&phi->instr);
nir_def_replace(&phi->def, merged);
}
static bool

View file

@ -184,8 +184,7 @@ lower_alu_instr(nir_builder *b, nir_alu_instr *alu, bool has_fcsel_ne,
if (rep) {
/* We've emitted a replacement instruction */
nir_def_rewrite_uses(&alu->def, rep);
nir_instr_remove(&alu->instr);
nir_def_replace(&alu->def, rep);
} else {
if (alu->def.bit_size == 1)
alu->def.bit_size = 32;

View file

@ -402,8 +402,7 @@ nir_lower_const_arrays_to_uniforms(nir_shader *shader,
nir_def *new_def = nir_load_deref(&b, new_deref_instr);
nir_def_rewrite_uses(&intrin->def, new_def);
nir_instr_remove(&intrin->instr);
nir_def_replace(&intrin->def, new_def);
}
}

View file

@ -59,8 +59,7 @@ lower_input_io(nir_builder *b, nir_intrinsic_instr *intr, void *data)
intr->def.bit_size, intr->src[1].ssa);
nir_intrinsic_instr *new_intr = nir_instr_as_intrinsic(load->parent_instr);
nir_intrinsic_copy_const_indices(new_intr, intr);
nir_def_rewrite_uses(&intr->def, load);
nir_instr_remove(&intr->instr);
nir_def_replace(&intr->def, load);
return true;
}
bool

View file

@ -56,8 +56,7 @@ lower_cube_size(nir_builder *b, nir_intrinsic_instr *intrin)
}
nir_def *vec = nir_vec_scalars(b, comps, intrin->def.num_components);
nir_def_rewrite_uses(&intrin->def, vec);
nir_instr_remove(&intrin->instr);
nir_def_replace(&intrin->def, vec);
nir_instr_free(&intrin->instr);
}
@ -154,9 +153,7 @@ lower_image_samples_identical_to_fragment_mask_load(nir_builder *b, nir_intrinsi
nir_builder_instr_insert(b, &fmask_load->instr);
nir_def *samples_identical = nir_ieq_imm(b, &fmask_load->def, 0);
nir_def_rewrite_uses(&intrin->def, samples_identical);
nir_instr_remove(&intrin->instr);
nir_def_replace(&intrin->def, samples_identical);
nir_instr_free(&intrin->instr);
}

View file

@ -247,8 +247,7 @@ lower_alu_instr(nir_builder *b, nir_alu_instr *alu)
if (rep) {
/* We've emitted a replacement instruction */
nir_def_rewrite_uses(&alu->def, rep);
nir_instr_remove(&alu->instr);
nir_def_replace(&alu->def, rep);
}
return true;

View file

@ -2276,8 +2276,7 @@ lower_explicit_io_array_length(nir_builder *b, nir_intrinsic_instr *intrin,
nir_def *remaining = nir_usub_sat(b, size, offset);
nir_def *arr_size = nir_udiv_imm(b, remaining, stride);
nir_def_rewrite_uses(&intrin->def, arr_size);
nir_instr_remove(&intrin->instr);
nir_def_replace(&intrin->def, arr_size);
}
static void
@ -3115,8 +3114,7 @@ nir_lower_color_inputs(nir_shader *nir)
load = nir_channels(&b, load, BITFIELD_RANGE(start, count));
}
nir_def_rewrite_uses(&intrin->def, load);
nir_instr_remove(instr);
nir_def_replace(&intrin->def, load);
progress = true;
}
}

View file

@ -77,9 +77,7 @@ lower_load_input_to_scalar(nir_builder *b, nir_intrinsic_instr *intr)
loads[i] = &chan_intr->def;
}
nir_def_rewrite_uses(&intr->def,
nir_vec(b, loads, intr->num_components));
nir_instr_remove(&intr->instr);
nir_def_replace(&intr->def, nir_vec(b, loads, intr->num_components));
}
static void
@ -124,9 +122,7 @@ lower_load_to_scalar(nir_builder *b, nir_intrinsic_instr *intr)
loads[i] = &chan_intr->def;
}
nir_def_rewrite_uses(&intr->def,
nir_vec(b, loads, intr->num_components));
nir_instr_remove(&intr->instr);
nir_def_replace(&intr->def, nir_vec(b, loads, intr->num_components));
}
static void
@ -432,11 +428,7 @@ lower_load_to_scalar_early(nir_builder *b, nir_intrinsic_instr *intr,
loads[i] = &chan_intr->def;
}
nir_def_rewrite_uses(&intr->def,
nir_vec(b, loads, intr->num_components));
/* Remove the old load intrinsic */
nir_instr_remove(&intr->instr);
nir_def_replace(&intr->def, nir_vec(b, loads, intr->num_components));
}
static void

View file

@ -244,8 +244,7 @@ fixup_interpolation_instr(struct lower_io_state *state,
* correct part of the temporary.
*/
nir_def *load = nir_load_deref(b, nir_src_as_deref(interp->src[0]));
nir_def_rewrite_uses(&interp->def, load);
nir_instr_remove(&interp->instr);
nir_def_replace(&interp->def, load);
nir_deref_path_finish(&interp_path);
}

View file

@ -61,8 +61,7 @@ lower_load_and_store_is_helper(nir_builder *b,
case nir_intrinsic_is_helper_invocation: {
b->cursor = nir_before_instr(&intrin->instr);
nir_def *is_helper = nir_load_deref(b, is_helper_deref);
nir_def_rewrite_uses(&intrin->def, is_helper);
nir_instr_remove(&intrin->instr);
nir_def_replace(&intrin->def, is_helper);
return true;
}
default:

View file

@ -57,8 +57,7 @@ lower_load_const_instr_scalar(nir_load_const_instr *lower)
nir_def *vec = nir_vec(&b, loads, lower->def.num_components);
/* Replace the old load with a reference to our reconstructed vector. */
nir_def_rewrite_uses(&lower->def, vec);
nir_instr_remove(&lower->instr);
nir_def_replace(&lower->def, vec);
return true;
}

View file

@ -217,8 +217,7 @@ lower_mem_load(nir_builder *b, nir_intrinsic_instr *intrin,
nir_def *result = nir_extract_bits(b, chunks, num_chunks, 0,
num_components, bit_size);
nir_def_rewrite_uses(&intrin->def, result);
nir_instr_remove(&intrin->instr);
nir_def_replace(&intrin->def, result);
return true;
}

View file

@ -179,8 +179,7 @@ lower_pack_instr(nir_builder *b, nir_instr *instr, void *data)
default:
unreachable("Impossible opcode");
}
nir_def_rewrite_uses(&alu_instr->def, dest);
nir_instr_remove(&alu_instr->instr);
nir_def_replace(&alu_instr->def, dest);
return true;
}

View file

@ -84,9 +84,7 @@ nir_lower_patch_vertices(nir_shader *nir,
}
progress = true;
nir_def_rewrite_uses(&intr->def,
val);
nir_instr_remove(instr);
nir_def_replace(&intr->def, val);
}
}
}

View file

@ -234,10 +234,7 @@ lower_phis_to_scalar_block(nir_block *block,
nir_instr_insert_after(&last_phi->instr, &vec->instr);
nir_def_rewrite_uses(&phi->def,
&vec->def);
nir_instr_remove(&phi->instr);
nir_def_replace(&phi->def, &vec->def);
exec_list_push_tail(&state->dead_instrs, &phi->instr.node);
progress = true;

View file

@ -119,8 +119,7 @@ lower_printf_intrin(nir_builder *b, nir_intrinsic_instr *prntf, void *_options)
nir_pop_if(b, NULL);
nir_def *ret_val = nir_if_phi(b, printf_succ_val, printf_fail_val);
nir_def_rewrite_uses(&prntf->def, ret_val);
nir_instr_remove(&prntf->instr);
nir_def_replace(&prntf->def, ret_val);
return true;
}

View file

@ -172,8 +172,7 @@ lower_readonly_image_instr_intrinsic(nir_builder *b, nir_intrinsic_instr *intrin
nir_def *res = nir_trim_vector(b, &tex->def,
intrin->def.num_components);
nir_def_rewrite_uses(&intrin->def, res);
nir_instr_remove(&intrin->instr);
nir_def_replace(&intrin->def, res);
return true;
}

View file

@ -71,8 +71,7 @@ rewrite_load(nir_intrinsic_instr *load, struct regs_to_ssa_state *state)
nir_intrinsic_instr *decl = nir_instr_as_intrinsic(reg->parent_instr);
nir_def *def = nir_phi_builder_value_get_block_def(value, block);
nir_def_rewrite_uses(&load->def, def);
nir_instr_remove(&load->instr);
nir_def_replace(&load->def, def);
if (nir_def_is_unused(&decl->def))
nir_instr_remove(&decl->instr);

View file

@ -78,8 +78,7 @@ lower_single_sampled_instr(nir_builder *b,
return false;
}
nir_def_rewrite_uses(&intrin->def, lowered);
nir_instr_remove(&intrin->instr);
nir_def_replace(&intrin->def, lowered);
return true;
}

View file

@ -353,8 +353,7 @@ lower_zero_lod(nir_builder *b, nir_tex_instr *tex)
b->cursor = nir_before_instr(&tex->instr);
if (tex->op == nir_texop_lod) {
nir_def_rewrite_uses(&tex->def, nir_imm_int(b, 0));
nir_instr_remove(&tex->instr);
nir_def_replace(&tex->def, nir_imm_int(b, 0));
return;
}
@ -959,8 +958,7 @@ lower_tex_to_txd(nir_builder *b, nir_tex_instr *tex)
tex->def.num_components,
tex->def.bit_size);
nir_builder_instr_insert(b, &txd->instr);
nir_def_rewrite_uses(&tex->def, &txd->def);
nir_instr_remove(&tex->instr);
nir_def_replace(&tex->def, &txd->def);
return txd;
}
@ -999,8 +997,7 @@ lower_txb_to_txl(nir_builder *b, nir_tex_instr *tex)
tex->def.num_components,
tex->def.bit_size);
nir_builder_instr_insert(b, &txl->instr);
nir_def_rewrite_uses(&tex->def, &txl->def);
nir_instr_remove(&tex->instr);
nir_def_replace(&tex->def, &txl->def);
return txl;
}
@ -1303,8 +1300,7 @@ lower_tg4_offsets(nir_builder *b, nir_tex_instr *tex)
dest[4] = nir_get_scalar(residency, 0);
nir_def *res = nir_vec_scalars(b, dest, tex->def.num_components);
nir_def_rewrite_uses(&tex->def, res);
nir_instr_remove(&tex->instr);
nir_def_replace(&tex->def, res);
return true;
}

View file

@ -108,9 +108,7 @@ nir_lower_uniforms_to_ubo_instr(nir_builder *b, nir_instr *instr, void *data)
nir_intrinsic_set_range_base(load, nir_intrinsic_base(intr) * multiplier);
nir_intrinsic_set_range(load, nir_intrinsic_range(intr) * multiplier);
}
nir_def_rewrite_uses(&intr->def, load_result);
nir_instr_remove(&intr->instr);
nir_def_replace(&intr->def, load_result);
return true;
}

View file

@ -668,8 +668,7 @@ rename_variables(struct lower_variables_state *state)
val = nir_vector_extract(&b, val, comp);
}
nir_def_rewrite_uses(&intrin->def, val);
nir_instr_remove(&intrin->instr);
nir_def_replace(&intrin->def, val);
break;
}

View file

@ -87,8 +87,7 @@ try_fold_alu(nir_builder *b, nir_alu_instr *alu)
nir_def *imm = nir_build_imm(b, alu->def.num_components,
alu->def.bit_size,
dest);
nir_def_rewrite_uses(&alu->def, imm);
nir_instr_remove(&alu->instr);
nir_def_replace(&alu->def, imm);
nir_instr_free(&alu->instr);
return true;
@ -206,8 +205,7 @@ try_fold_intrinsic(nir_builder *b, nir_intrinsic_instr *intrin,
b->cursor = nir_before_instr(&intrin->instr);
nir_def *val = nir_build_imm(b, intrin->def.num_components,
intrin->def.bit_size, v);
nir_def_rewrite_uses(&intrin->def, val);
nir_instr_remove(&intrin->instr);
nir_def_replace(&intrin->def, val);
return true;
}
return false;
@ -245,8 +243,7 @@ try_fold_intrinsic(nir_builder *b, nir_intrinsic_instr *intrin,
val = nir_build_imm(b, intrin->def.num_components,
intrin->def.bit_size, imm);
}
nir_def_rewrite_uses(&intrin->def, val);
nir_instr_remove(&intrin->instr);
nir_def_replace(&intrin->def, val);
return true;
}
@ -270,9 +267,7 @@ try_fold_intrinsic(nir_builder *b, nir_intrinsic_instr *intrin,
* the data is constant.
*/
if (nir_src_is_const(intrin->src[0])) {
nir_def_rewrite_uses(&intrin->def,
intrin->src[0].ssa);
nir_instr_remove(&intrin->instr);
nir_def_replace(&intrin->def, intrin->src[0].ssa);
return true;
}
return false;
@ -281,9 +276,7 @@ try_fold_intrinsic(nir_builder *b, nir_intrinsic_instr *intrin,
case nir_intrinsic_vote_ieq:
if (nir_src_is_const(intrin->src[0])) {
b->cursor = nir_before_instr(&intrin->instr);
nir_def_rewrite_uses(&intrin->def,
nir_imm_true(b));
nir_instr_remove(&intrin->instr);
nir_def_replace(&intrin->def, nir_imm_true(b));
return true;
}
return false;

View file

@ -97,8 +97,7 @@ opt_constant_if(nir_if *if_stmt, bool condition)
}
assert(def);
nir_def_rewrite_uses(&phi->def, def);
nir_instr_remove(&phi->instr);
nir_def_replace(&phi->def, def);
}
}

View file

@ -213,8 +213,7 @@ nir_opt_idiv_const_instr(nir_builder *b, nir_instr *instr, void *user_data)
}
nir_def *qvec = nir_vec(b, q, alu->def.num_components);
nir_def_rewrite_uses(&alu->def, qvec);
nir_instr_remove(&alu->instr);
nir_def_replace(&alu->def, qvec);
return true;
}

View file

@ -233,9 +233,7 @@ opt_intrinsics_alu(nir_builder *b, nir_alu_instr *alu,
}
if (replacement) {
nir_def_rewrite_uses(&alu->def,
replacement);
nir_instr_remove(&alu->instr);
nir_def_replace(&alu->def, replacement);
return true;
} else {
return false;
@ -297,8 +295,7 @@ try_opt_exclusive_scan_to_inclusive(nir_intrinsic_instr *intrin)
nir_foreach_use_including_if_safe(src, &intrin->def) {
/* Remove alu. */
nir_alu_instr *alu = nir_instr_as_alu(nir_src_parent_instr(src));
nir_def_rewrite_uses(&alu->def, &intrin->def);
nir_instr_remove(&alu->instr);
nir_def_replace(&alu->def, &intrin->def);
}
return true;
@ -337,9 +334,7 @@ opt_intrinsics_intrin(nir_builder *b, nir_intrinsic_instr *intrin,
if (alu->op == nir_op_ine)
new_expr = nir_inot(b, new_expr);
nir_def_rewrite_uses(&alu->def,
new_expr);
nir_instr_remove(&alu->instr);
nir_def_replace(&alu->def, new_expr);
progress = true;
}
}

View file

@ -548,15 +548,12 @@ nir_opt_large_constants(nir_shader *shader,
if (info->is_small) {
b.cursor = nir_after_instr(&intrin->instr);
nir_def *val = build_small_constant_load(&b, deref, info, size_align);
nir_def_rewrite_uses(&intrin->def, val);
nir_instr_remove(&intrin->instr);
nir_def_replace(&intrin->def, val);
nir_deref_instr_remove_if_unused(deref);
} else if (info->is_constant) {
b.cursor = nir_after_instr(&intrin->instr);
nir_def *val = build_constant_load(&b, deref, size_align);
nir_def_rewrite_uses(&intrin->def,
val);
nir_instr_remove(&intrin->instr);
nir_def_replace(&intrin->def, val);
nir_deref_instr_remove_if_unused(deref);
}
break;

View file

@ -124,8 +124,7 @@ remove_phis_block(nir_block *block, nir_builder *b)
def = nir_mov_alu(b, mov->src[0], def->num_components);
}
nir_def_rewrite_uses(&phi->def, def);
nir_instr_remove(&phi->instr);
nir_def_replace(&phi->def, def);
progress = true;
}

View file

@ -1796,8 +1796,7 @@ remove_dead_varyings(struct linkage_info *linkage,
else
replacement = nir_undef(b, 1, loadi->def.bit_size);
nir_def_rewrite_uses(&loadi->def, replacement);
nir_instr_remove(&loadi->instr);
nir_def_replace(&loadi->def, replacement);
*progress |= list_index ? nir_progress_producer :
nir_progress_consumer;
@ -2189,8 +2188,7 @@ propagate_uniform_expressions(struct linkage_info *linkage,
clone = build_convert_inf_to_nan(b, clone);
/* Replace the original load. */
nir_def_rewrite_uses(&loadi->def, clone);
nir_instr_remove(&loadi->instr);
nir_def_replace(&loadi->def, clone);
*progress |= list_index ? nir_progress_producer :
nir_progress_consumer;
}

View file

@ -152,8 +152,7 @@ vectorize_load(nir_intrinsic_instr *chan[8], unsigned start, unsigned count,
}
} else {
for (unsigned i = start; i < start + count; i++) {
nir_def_rewrite_uses(&chan[i]->def, nir_channel(&b, def, i - start));
nir_instr_remove(&chan[i]->instr);
nir_def_replace(&chan[i]->def, nir_channel(&b, def, i - start));
}
}
}

View file

@ -352,9 +352,7 @@ lower_ubo_load_to_uniform(nir_intrinsic_instr *instr, nir_builder *b,
nir_load_uniform(b, instr->num_components, instr->def.bit_size,
uniform_offset, .base = const_offset);
nir_def_rewrite_uses(&instr->def, uniform);
nir_instr_remove(&instr->instr);
nir_def_replace(&instr->def, uniform);
return true;
}

View file

@ -348,8 +348,7 @@ lower_block_to_explicit_input(nir_block *block, nir_builder *b,
b->cursor = nir_before_instr(&intr->instr);
nir_def *iid = build_invocation_id(b, state);
nir_def_rewrite_uses(&intr->def, iid);
nir_instr_remove(&intr->instr);
nir_def_replace(&intr->def, iid);
break;
}

View file

@ -138,9 +138,7 @@ lower_load_push_constant(struct tu_device *dev,
nir_ushr_imm(b, instr->src[0].ssa, 2),
.base = base);
nir_def_rewrite_uses(&instr->def, load);
nir_instr_remove(&instr->instr);
nir_def_replace(&instr->def, load);
}
static void
@ -216,8 +214,7 @@ lower_vulkan_resource_index(struct tu_device *dev, nir_builder *b,
nir_ishl(b, vulkan_idx, shift)),
shift);
nir_def_rewrite_uses(&instr->def, def);
nir_instr_remove(&instr->instr);
nir_def_replace(&instr->def, def);
}
static void
@ -233,8 +230,7 @@ lower_vulkan_resource_reindex(nir_builder *b, nir_intrinsic_instr *instr)
nir_ishl(b, delta, shift)),
shift);
nir_def_rewrite_uses(&instr->def, new_index);
nir_instr_remove(&instr->instr);
nir_def_replace(&instr->def, new_index);
}
static void
@ -248,8 +244,7 @@ lower_load_vulkan_descriptor(nir_builder *b, nir_intrinsic_instr *intrin)
nir_vec3(b, nir_channel(b, old_index, 0),
nir_channel(b, old_index, 1),
nir_imm_int(b, 0));
nir_def_rewrite_uses(&intrin->def, new_index);
nir_instr_remove(&intrin->instr);
nir_def_replace(&intrin->def, new_index);
}
static bool
@ -480,8 +475,7 @@ lower_intrinsic(nir_builder *b, nir_intrinsic_instr *instr,
ir3_load_driver_ubo_indirect(b, 2, &shader->const_state.fdm_ubo,
param, view, nir_intrinsic_range(instr));
nir_def_rewrite_uses(&instr->def, result);
nir_instr_remove(&instr->instr);
nir_def_replace(&instr->def, result);
return true;
}
case nir_intrinsic_load_frag_invocation_count: {
@ -492,8 +486,7 @@ lower_intrinsic(nir_builder *b, nir_intrinsic_instr *instr,
ir3_load_driver_ubo(b, 1, &shader->const_state.fdm_ubo,
IR3_DP_FS_FRAG_INVOCATION_COUNT);
nir_def_rewrite_uses(&instr->def, result);
nir_instr_remove(&instr->instr);
nir_def_replace(&instr->def, result);
return true;
}
@ -717,8 +710,7 @@ lower_inline_ubo(nir_builder *b, nir_intrinsic_instr *intrin, void *cb_data)
nir_ishr_imm(b, offset, 2), .base = base);
}
nir_def_rewrite_uses(&intrin->def, val);
nir_instr_remove(&intrin->instr);
nir_def_replace(&intrin->def, val);
return true;
}

View file

@ -3547,8 +3547,7 @@ nir_to_tgsi_lower_64bit_load_const(nir_builder *b, nir_load_const_instr *instr)
num_components == 4 ? nir_channel(b, &second->def, 1) : NULL,
};
nir_def *new = nir_vec(b, channels, num_components);
nir_def_rewrite_uses(&instr->def, new);
nir_instr_remove(&instr->instr);
nir_def_replace(&instr->def, new);
return true;
}

View file

@ -497,9 +497,7 @@ crocus_setup_uniforms(ASSERTED const struct intel_device_info *devinfo,
intrin->def.bit_size);
nir_builder_instr_insert(&b, &load_ubo->instr);
nir_def_rewrite_uses(&intrin->def,
&load_ubo->def);
nir_instr_remove(&intrin->instr);
nir_def_replace(&intrin->def, &load_ubo->def);
continue;
}
case nir_intrinsic_load_user_clip_plane: {

View file

@ -268,8 +268,7 @@ lower_load_draw_params(nir_builder *b, nir_intrinsic_instr *intr,
unsigned channel = intr->intrinsic == nir_intrinsic_load_first_vertex ? 0 :
intr->intrinsic == nir_intrinsic_load_base_instance ? 1 :
intr->intrinsic == nir_intrinsic_load_draw_id ? 2 : 3;
nir_def_rewrite_uses(&intr->def, nir_channel(b, load, channel));
nir_instr_remove(&intr->instr);
nir_def_replace(&intr->def, nir_channel(b, load, channel));
return true;
}
@ -297,8 +296,7 @@ lower_load_patch_vertices_in(nir_builder *b, nir_intrinsic_instr *intr,
nir_def *load = b->shader->info.stage == MESA_SHADER_TESS_CTRL ?
d3d12_get_state_var(b, D3D12_STATE_VAR_PATCH_VERTICES_IN, "d3d12_FirstVertex", glsl_uint_type(), _state) :
nir_imm_int(b, b->shader->info.tess.tcs_vertices_out);
nir_def_rewrite_uses(&intr->def, load);
nir_instr_remove(&intr->instr);
nir_def_replace(&intr->def, load);
return true;
}
@ -482,10 +480,7 @@ lower_instr(nir_intrinsic_instr *instr, nir_builder *b,
.range = ~0,
);
nir_def_rewrite_uses(&instr->def, load);
/* Remove the old load_* instruction and any parent derefs */
nir_instr_remove(&instr->instr);
nir_def_replace(&instr->def, load);
for (nir_deref_instr *d = deref; d; d = nir_deref_instr_parent(d)) {
/* If anyone is using this deref, leave it alone */
if (!list_is_empty(&d->def.uses))

View file

@ -880,8 +880,7 @@ lower_alu(struct etna_compile *c, nir_alu_instr *alu)
nir_def *def = nir_build_imm(&b, num_components, 32, value);
if (num_components == info->num_inputs) {
nir_def_rewrite_uses(&alu->def, def);
nir_instr_remove(&alu->instr);
nir_def_replace(&alu->def, def);
return;
}

View file

@ -51,9 +51,7 @@ lower_load_uniform_to_scalar(nir_builder *b, nir_intrinsic_instr *intr)
loads[i] = &chan_intr->def;
}
nir_def_rewrite_uses(&intr->def,
nir_vec(b, loads, intr->num_components));
nir_instr_remove(&intr->instr);
nir_def_replace(&intr->def, nir_vec(b, loads, intr->num_components));
}
void

View file

@ -80,9 +80,7 @@ lima_nir_split_load_input_instr(nir_builder *b,
new_intrin->src[0] = nir_src_for_ssa(intrin->src[0].ssa);
nir_builder_instr_insert(b, &new_intrin->instr);
nir_def_rewrite_uses(&alu->def,
&new_intrin->def);
nir_instr_remove(&alu->instr);
nir_def_replace(&alu->def, &new_intrin->def);
return true;
}

View file

@ -85,9 +85,7 @@ lower_load_poly_line_smooth_enabled(nir_shader *nir,
continue;
b.cursor = nir_before_instr(instr);
nir_def_rewrite_uses(&intrin->def, nir_imm_true(&b));
nir_instr_remove(instr);
nir_def_replace(&intrin->def, nir_imm_true(&b));
nir_instr_free(instr);
}
}

View file

@ -117,8 +117,7 @@ r300_nir_lower_fcsel_instr(nir_builder *b, nir_instr *instr, void *data)
nir_ssa_for_alu_src(b, alu, 1), slt);
}
nir_def_rewrite_uses(&alu->def, lrp);
nir_instr_remove(&alu->instr);
nir_def_replace(&alu->def, lrp);
return true;
}
return false;

View file

@ -1405,8 +1405,7 @@ r600_lower_64bit_load_const(nir_builder *b, nir_load_const_instr *instr)
num_components == 4 ? nir_channel(b, &second->def, 1) : NULL,
};
nir_def *new_ir = nir_vec(b, channels, num_components);
nir_def_rewrite_uses(&instr->def, new_ir);
nir_instr_remove(&instr->instr);
nir_def_replace(&instr->def, new_ir);
return true;
}

View file

@ -325,8 +325,7 @@ r600_lower_tess_io_impl(nir_builder *b, nir_instr *instr, enum mesa_prim prim_ty
auto base = nir_load_tcs_in_param_base_r600(b);
vertices_in = nir_channel(b, base, 2);
}
nir_def_rewrite_uses(&op->def, vertices_in);
nir_instr_remove(&op->instr);
nir_def_replace(&op->def, vertices_in);
return true;
}
case nir_intrinsic_load_per_vertex_input: {

View file

@ -147,10 +147,7 @@ r600_create_new_load(nir_builder *b,
for (unsigned i = 0; i < old_num_comps; ++i)
channels[i] = comp - var->data.location_frac + i;
nir_def *load = nir_swizzle(b, &new_intr->def, channels, old_num_comps);
nir_def_rewrite_uses(&intr->def, load);
/* Remove the old load intrinsic */
nir_instr_remove(&intr->instr);
nir_def_replace(&intr->def, load);
}
static bool

View file

@ -303,8 +303,7 @@ static bool lower_resource_intrinsic(nir_builder *b, nir_intrinsic_instr *intrin
nir_def *desc = load_ssbo_desc(b, &intrin->src[0], s);
nir_def *size = nir_channel(b, desc, 2);
nir_def_rewrite_uses(&intrin->def, size);
nir_instr_remove(&intrin->instr);
nir_def_replace(&intrin->def, size);
break;
}
case nir_intrinsic_image_deref_load:
@ -335,8 +334,7 @@ static bool lower_resource_intrinsic(nir_builder *b, nir_intrinsic_instr *intrin
nir_def *desc = load_deref_image_desc(b, deref, desc_type, is_load, s);
if (intrin->intrinsic == nir_intrinsic_image_deref_descriptor_amd) {
nir_def_rewrite_uses(&intrin->def, desc);
nir_instr_remove(&intrin->instr);
nir_def_replace(&intrin->def, desc);
} else {
nir_intrinsic_set_image_dim(intrin, glsl_get_sampler_dim(deref->type));
nir_intrinsic_set_image_array(intrin, glsl_sampler_type_is_array(deref->type));
@ -377,8 +375,7 @@ static bool lower_resource_intrinsic(nir_builder *b, nir_intrinsic_instr *intrin
nir_def *desc = load_bindless_image_desc(b, index, desc_type, is_load, s);
if (intrin->intrinsic == nir_intrinsic_bindless_image_descriptor_amd) {
nir_def_rewrite_uses(&intrin->def, desc);
nir_instr_remove(&intrin->instr);
nir_def_replace(&intrin->def, desc);
} else {
nir_src_rewrite(&intrin->src[0], desc);
}
@ -516,8 +513,7 @@ static bool lower_resource_tex(nir_builder *b, nir_tex_instr *tex,
image = load_deref_sampler_desc(b, texture_deref, desc_type, s, true);
else
image = load_bindless_sampler_desc(b, texture_handle, desc_type, s);
nir_def_rewrite_uses(&tex->def, image);
nir_instr_remove(&tex->instr);
nir_def_replace(&tex->def, image);
return true;
}
@ -527,8 +523,7 @@ static bool lower_resource_tex(nir_builder *b, nir_tex_instr *tex,
sampler = load_deref_sampler_desc(b, sampler_deref, AC_DESC_SAMPLER, s, true);
else
sampler = load_bindless_sampler_desc(b, sampler_handle, AC_DESC_SAMPLER, s);
nir_def_rewrite_uses(&tex->def, sampler);
nir_instr_remove(&tex->instr);
nir_def_replace(&tex->def, sampler);
return true;
}

View file

@ -578,8 +578,7 @@ lower_vs_input_instr(nir_builder *b, nir_intrinsic_instr *intrin, void *state)
nir_def *replacement = nir_vec(b, &comp[component], num_components);
nir_def_rewrite_uses(&intrin->def, replacement);
nir_instr_remove(&intrin->instr);
nir_def_replace(&intrin->def, replacement);
nir_instr_free(&intrin->instr);
return true;

View file

@ -2104,9 +2104,7 @@ static bool lower_ps_load_color_intrinsic(nir_builder *b, nir_instr *instr, void
unsigned index = intrin->intrinsic == nir_intrinsic_load_color0 ? 0 : 1;
assert(colors[index]);
nir_def_rewrite_uses(&intrin->def, colors[index]);
nir_instr_remove(&intrin->instr);
nir_def_replace(&intrin->def, colors[index]);
return true;
}

View file

@ -49,8 +49,7 @@ replace_intrinsic_with_vec(nir_builder *b, nir_intrinsic_instr *intr,
/* Replace the old intrinsic with a reference to our reconstructed
* vector.
*/
nir_def_rewrite_uses(&intr->def, vec);
nir_instr_remove(&intr->instr);
nir_def_replace(&intr->def, vec);
}
static nir_def *

View file

@ -1057,8 +1057,7 @@ lower_64bit_pack_instr(nir_builder *b, nir_instr *instr, void *data)
default:
unreachable("Impossible opcode");
}
nir_def_rewrite_uses(&alu_instr->def, dest);
nir_instr_remove(&alu_instr->instr);
nir_def_replace(&alu_instr->def, dest);
return true;
}
@ -1197,8 +1196,7 @@ lower_system_values_to_inlined_uniforms_instr(nir_builder *b,
new_dest_def = dwords[0];
else
new_dest_def = nir_pack_64_2x32_split(b, dwords[0], dwords[1]);
nir_def_rewrite_uses(&intrin->def, new_dest_def);
nir_instr_remove(&intrin->instr);
nir_def_replace(&intrin->def, new_dest_def);
return true;
}
@ -2370,8 +2368,7 @@ rewrite_atomic_ssbo_instr(nir_builder *b, nir_instr *instr, struct bo_vars *bo)
}
nir_def *load = nir_vec(b, result, num_components);
nir_def_rewrite_uses(&intr->def, load);
nir_instr_remove(instr);
nir_def_replace(&intr->def, load);
}
static bool
@ -2804,8 +2801,7 @@ rewrite_read_as_0(nir_builder *b, nir_instr *instr, void *data)
break;
}
}
nir_def_rewrite_uses(&intr->def, zero);
nir_instr_remove(instr);
nir_def_replace(&intr->def, zero);
return true;
}

View file

@ -204,8 +204,7 @@ lvp_inline_uniforms(nir_shader *nir, const struct lvp_shader *shader, const uint
if (offset == uniform_dw_offsets[i]) {
b.cursor = nir_before_instr(&intr->instr);
nir_def *def = nir_imm_int(&b, uniform_values[i]);
nir_def_rewrite_uses(&intr->def, def);
nir_instr_remove(&intr->instr);
nir_def_replace(&intr->def, def);
break;
}
}
@ -248,9 +247,8 @@ lvp_inline_uniforms(nir_shader *nir, const struct lvp_shader *shader, const uint
}
/* Replace the original uniform load. */
nir_def_rewrite_uses(&intr->def,
nir_vec(&b, components, num_components));
nir_instr_remove(&intr->instr);
nir_def_replace(&intr->def,
nir_vec(&b, components, num_components));
}
}
}

View file

@ -110,12 +110,10 @@ lvp_lower_node_payload_intrinsic(nir_builder *b, nir_intrinsic_instr *intr,
nir_instr_remove(&intr->instr);
return true;
case nir_intrinsic_finalize_incoming_node_payload:
nir_def_rewrite_uses(&intr->def, nir_imm_true(b));
nir_instr_remove(&intr->instr);
nir_def_replace(&intr->def, nir_imm_true(b));
return true;
case nir_intrinsic_load_coalesced_input_count:
nir_def_rewrite_uses(&intr->def, nir_imm_int(b, 1));
nir_instr_remove(&intr->instr);
nir_def_replace(&intr->def, nir_imm_int(b, 1));
return true;
default:
return false;

View file

@ -128,8 +128,7 @@ lvp_lower_ray_tracing_derefs(nir_shader *shader)
nir_def *offset = is_shader_call_data ? arg_offset : nir_imm_int(b, 0);
nir_deref_instr *replacement =
nir_build_deref_cast(b, offset, nir_var_function_temp, deref->var->type, 0);
nir_def_rewrite_uses(&deref->def, &replacement->def);
nir_instr_remove(&deref->instr);
nir_def_replace(&deref->def, &replacement->def);
}
}
}
@ -542,8 +541,7 @@ lvp_lower_isec_intrinsic(nir_builder *b, nir_intrinsic_instr *instr, void *data)
}
nir_pop_if(b, NULL);
nir_def_rewrite_uses(&instr->def, nir_load_var(b, commit));
nir_instr_remove(&instr->instr);
nir_def_replace(&instr->def, nir_load_var(b, commit));
return true;
}
@ -1026,8 +1024,7 @@ lvp_lower_ray_tracing_stack_base(nir_builder *b, nir_intrinsic_instr *instr, voi
b->cursor = nir_after_instr(&instr->instr);
nir_def_rewrite_uses(&instr->def, nir_imm_int(b, b->shader->scratch_size));
nir_instr_remove(&instr->instr);
nir_def_replace(&instr->def, nir_imm_int(b, b->shader->scratch_size));
return true;
}

View file

@ -49,8 +49,7 @@ static void lower_vulkan_resource_index(nir_builder *b,
nir_imm_int(b, desc_set),
nir_imm_int(b, binding),
nir_imm_int(b, desc_type));
nir_def_rewrite_uses(&intr->def, def);
nir_instr_remove(&intr->instr);
nir_def_replace(&intr->def, def);
}
static void lower_load_global_constant_to_scalar(nir_builder *b,
@ -84,9 +83,7 @@ static void lower_load_global_constant_to_scalar(nir_builder *b,
loads[i] = &chan_intr->def;
}
nir_def_rewrite_uses(&intr->def,
nir_vec(b, loads, intr->num_components));
nir_instr_remove(&intr->instr);
nir_def_replace(&intr->def, nir_vec(b, loads, intr->num_components));
}
static bool lower_intrinsic(nir_builder *b, nir_intrinsic_instr *instr)

View file

@ -543,11 +543,8 @@ nir_remove_llvm17_scratch(nir_shader *nir)
if (offset == NULL)
continue;
nir_def_rewrite_uses(&intrin->def,
rebuild_value_from_store(
&scratch_stores, &intrin->def,
nir_src_as_uint(intrin->src[0])));
nir_instr_remove(instr);
nir_def_replace(&intrin->def,
rebuild_value_from_store(&scratch_stores, &intrin->def, nir_src_as_uint(intrin->src[0])));
progress = true;
}

View file

@ -433,9 +433,7 @@ brw_nir_lower_vs_inputs(nir_shader *nir)
nir_def_init(&load->instr, &load->def, 1, 32);
nir_builder_instr_insert(&b, &load->instr);
nir_def_rewrite_uses(&intrin->def,
&load->def);
nir_instr_remove(&intrin->instr);
nir_def_replace(&intrin->def, &load->def);
break;
}
@ -546,8 +544,7 @@ lower_barycentric_per_sample(nir_builder *b,
nir_def *centroid =
nir_load_barycentric(b, nir_intrinsic_load_barycentric_sample,
nir_intrinsic_interp_mode(intrin));
nir_def_rewrite_uses(&intrin->def, centroid);
nir_instr_remove(&intrin->instr);
nir_def_replace(&intrin->def, centroid);
return true;
}
@ -1160,9 +1157,7 @@ brw_nir_zero_inputs_instr(struct nir_builder *b, nir_intrinsic_instr *intrin,
nir_def *zero = nir_imm_zero(b, 1, 32);
nir_def_rewrite_uses(&intrin->def, zero);
nir_instr_remove(&intrin->instr);
nir_def_replace(&intrin->def, zero);
return true;
}

View file

@ -272,8 +272,7 @@ lower_cs_intrinsics_convert_block(struct lower_intrinsics_state *state,
if (intrinsic->def.bit_size == 64)
sysval = nir_u2u64(b, sysval);
nir_def_rewrite_uses(&intrinsic->def, sysval);
nir_instr_remove(&intrinsic->instr);
nir_def_replace(&intrinsic->def, sysval);
state->progress = true;
}

View file

@ -92,15 +92,11 @@ lower_any_hit_for_intersection(nir_shader *any_hit)
break;
case nir_intrinsic_load_ray_t_max:
nir_def_rewrite_uses(&intrin->def,
hit_t);
nir_instr_remove(&intrin->instr);
nir_def_replace(&intrin->def, hit_t);
break;
case nir_intrinsic_load_ray_hit_kind:
nir_def_rewrite_uses(&intrin->def,
hit_kind);
nir_instr_remove(&intrin->instr);
nir_def_replace(&intrin->def, hit_kind);
break;
default:

View file

@ -340,9 +340,7 @@ lower_rt_intrinsics_impl(nir_function_impl *impl,
progress = true;
if (sysval) {
nir_def_rewrite_uses(&intrin->def,
sysval);
nir_instr_remove(&intrin->instr);
nir_def_replace(&intrin->def, sysval);
}
}
}

View file

@ -118,9 +118,7 @@ lower_rt_io_derefs(nir_shader *shader)
nir_build_deref_cast(&b, call_data_addr,
nir_var_function_temp,
deref->var->type, 0);
nir_def_rewrite_uses(&deref->def,
&cast->def);
nir_instr_remove(&deref->instr);
nir_def_replace(&deref->def, &cast->def);
progress = true;
}
} else if (nir_deref_mode_is(deref, nir_var_ray_hit_attrib)) {
@ -131,9 +129,7 @@ lower_rt_io_derefs(nir_shader *shader)
nir_build_deref_cast(&b, hit_attrib_addr,
nir_var_function_temp,
deref->type, 0);
nir_def_rewrite_uses(&deref->def,
&cast->def);
nir_instr_remove(&deref->instr);
nir_def_replace(&deref->def, &cast->def);
progress = true;
}
}
@ -520,9 +516,7 @@ brw_nir_create_raygen_trampoline(const struct brw_compiler *compiler,
b.cursor = nir_before_instr(&intrin->instr);
nir_def *global_arg_addr =
load_trampoline_param(&b, rt_disp_globals_addr, 1, 64);
nir_def_rewrite_uses(&intrin->def,
global_arg_addr);
nir_instr_remove(instr);
nir_def_replace(&intrin->def, global_arg_addr);
}
}

View file

@ -330,9 +330,7 @@ elk_nir_lower_vs_inputs(nir_shader *nir,
nir_def_init(&load->instr, &load->def, 1, 32);
nir_builder_instr_insert(&b, &load->instr);
nir_def_rewrite_uses(&intrin->def,
&load->def);
nir_instr_remove(&intrin->instr);
nir_def_replace(&intrin->def, &load->def);
break;
}
@ -451,8 +449,7 @@ lower_barycentric_per_sample(nir_builder *b,
nir_def *centroid =
nir_load_barycentric(b, nir_intrinsic_load_barycentric_sample,
nir_intrinsic_interp_mode(intrin));
nir_def_rewrite_uses(&intrin->def, centroid);
nir_instr_remove(&intrin->instr);
nir_def_replace(&intrin->def, centroid);
return true;
}
@ -1035,9 +1032,7 @@ elk_nir_zero_inputs_instr(struct nir_builder *b, nir_intrinsic_instr *intrin,
nir_def *zero = nir_imm_zero(b, 1, 32);
nir_def_rewrite_uses(&intrin->def, zero);
nir_instr_remove(&intrin->instr);
nir_def_replace(&intrin->def, zero);
return true;
}

View file

@ -272,8 +272,7 @@ lower_cs_intrinsics_convert_block(struct lower_intrinsics_state *state,
if (intrinsic->def.bit_size == 64)
sysval = nir_u2u64(b, sysval);
nir_def_rewrite_uses(&intrinsic->def, sysval);
nir_instr_remove(&intrinsic->instr);
nir_def_replace(&intrinsic->def, sysval);
state->progress = true;
}

View file

@ -32,8 +32,7 @@ split_conversion(nir_builder *b, nir_alu_instr *alu, nir_alu_type src_type,
nir_def *src = nir_ssa_for_alu_src(b, alu, 0);
nir_def *tmp = nir_type_convert(b, src, src_type, tmp_type, nir_rounding_mode_undef);
nir_def *res = nir_type_convert(b, tmp, tmp_type, dst_type, nir_rounding_mode_undef);
nir_def_rewrite_uses(&alu->def, res);
nir_instr_remove(&alu->instr);
nir_def_replace(&alu->def, res);
}
static bool

View file

@ -1449,8 +1449,7 @@ lower_load_accel_struct_desc(nir_builder *b,
assert(load_desc->def.bit_size == 64);
assert(load_desc->def.num_components == 1);
nir_def_rewrite_uses(&load_desc->def, desc);
nir_instr_remove(&load_desc->instr);
nir_def_replace(&load_desc->def, desc);
return true;
}
@ -1532,8 +1531,7 @@ lower_res_index_intrinsic(nir_builder *b, nir_intrinsic_instr *intrin,
assert(intrin->def.bit_size == index->bit_size);
assert(intrin->def.num_components == index->num_components);
nir_def_rewrite_uses(&intrin->def, index);
nir_instr_remove(&intrin->instr);
nir_def_replace(&intrin->def, index);
return true;
}
@ -1550,8 +1548,7 @@ lower_res_reindex_intrinsic(nir_builder *b, nir_intrinsic_instr *intrin,
assert(intrin->def.bit_size == index->bit_size);
assert(intrin->def.num_components == index->num_components);
nir_def_rewrite_uses(&intrin->def, index);
nir_instr_remove(&intrin->instr);
nir_def_replace(&intrin->def, index);
return true;
}
@ -1572,8 +1569,7 @@ lower_load_vulkan_descriptor(nir_builder *b, nir_intrinsic_instr *intrin,
assert(intrin->def.bit_size == desc->bit_size);
assert(intrin->def.num_components == desc->num_components);
nir_def_rewrite_uses(&intrin->def, desc);
nir_instr_remove(&intrin->instr);
nir_def_replace(&intrin->def, desc);
return true;
}
@ -1615,8 +1611,7 @@ lower_get_ssbo_size(nir_builder *b, nir_intrinsic_instr *intrin,
}
nir_def *size = nir_channel(b, desc_range, 2);
nir_def_rewrite_uses(&intrin->def, size);
nir_instr_remove(&intrin->instr);
nir_def_replace(&intrin->def, size);
return true;
}

View file

@ -275,13 +275,11 @@ anv_nir_lower_multiview(nir_shader *shader, uint32_t view_mask,
switch (load->intrinsic) {
case nir_intrinsic_load_instance_id:
if (&load->def != state.instance_id_with_views) {
nir_def_rewrite_uses(&load->def, build_instance_id(&state));
nir_instr_remove(&load->instr);
nir_def_replace(&load->def, build_instance_id(&state));
}
break;
case nir_intrinsic_load_view_index:
nir_def_rewrite_uses(&load->def, view_index);
nir_instr_remove(&load->instr);
nir_def_replace(&load->def, view_index);
break;
case nir_intrinsic_emit_vertex_with_counter:
/* In geometry shaders, outputs become undefined after every

View file

@ -105,8 +105,7 @@ lower_ubo_load_instr(nir_builder *b, nir_intrinsic_instr *load,
}
}
nir_def_rewrite_uses(&load->def, val);
nir_instr_remove(&load->instr);
nir_def_replace(&load->def, val);
return true;
}

View file

@ -705,8 +705,7 @@ lower_load_accel_struct_desc(nir_builder *b,
assert(load_desc->def.bit_size == 64);
assert(load_desc->def.num_components == 1);
nir_def_rewrite_uses(&load_desc->def, desc);
nir_instr_remove(&load_desc->instr);
nir_def_replace(&load_desc->def, desc);
return true;
}
@ -755,8 +754,7 @@ lower_res_index_intrinsic(nir_builder *b, nir_intrinsic_instr *intrin,
assert(intrin->def.bit_size == index->bit_size);
assert(intrin->def.num_components == index->num_components);
nir_def_rewrite_uses(&intrin->def, index);
nir_instr_remove(&intrin->instr);
nir_def_replace(&intrin->def, index);
return true;
}
@ -777,8 +775,7 @@ lower_res_reindex_intrinsic(nir_builder *b, nir_intrinsic_instr *intrin,
assert(intrin->def.bit_size == index->bit_size);
assert(intrin->def.num_components == index->num_components);
nir_def_rewrite_uses(&intrin->def, index);
nir_instr_remove(&intrin->instr);
nir_def_replace(&intrin->def, index);
return true;
}
@ -798,8 +795,7 @@ lower_load_vulkan_descriptor(nir_builder *b, nir_intrinsic_instr *intrin,
assert(intrin->def.bit_size == desc->bit_size);
assert(intrin->def.num_components == desc->num_components);
nir_def_rewrite_uses(&intrin->def, desc);
nir_instr_remove(&intrin->instr);
nir_def_replace(&intrin->def, desc);
return true;
}
@ -824,8 +820,7 @@ lower_get_ssbo_size(nir_builder *b, nir_intrinsic_instr *intrin,
case nir_address_format_64bit_global_32bit_offset:
case nir_address_format_64bit_bounded_global: {
nir_def *size = nir_channel(b, desc, 2);
nir_def_rewrite_uses(&intrin->def, size);
nir_instr_remove(&intrin->instr);
nir_def_replace(&intrin->def, size);
break;
}

View file

@ -250,13 +250,11 @@ anv_nir_lower_multiview(nir_shader *shader, uint32_t view_mask)
switch (load->intrinsic) {
case nir_intrinsic_load_instance_id:
if (&load->def != state.instance_id_with_views) {
nir_def_rewrite_uses(&load->def, build_instance_id(&state));
nir_instr_remove(&load->instr);
nir_def_replace(&load->def, build_instance_id(&state));
}
break;
case nir_intrinsic_load_view_index:
nir_def_rewrite_uses(&load->def, view_index);
nir_instr_remove(&load->instr);
nir_def_replace(&load->def, view_index);
break;
case nir_intrinsic_emit_vertex_with_counter:
/* In geometry shaders, outputs become undefined after every

View file

@ -105,8 +105,7 @@ lower_ubo_load_instr(nir_builder *b, nir_intrinsic_instr *load,
}
}
nir_def_rewrite_uses(&load->def, val);
nir_instr_remove(&load->instr);
nir_def_replace(&load->def, val);
return true;
}

View file

@ -327,8 +327,7 @@ anv_nir_lower_ycbcr_textures_instr(nir_builder *builder,
swizzled_bpcs);
}
nir_def_rewrite_uses(&tex->def, result);
nir_instr_remove(&tex->instr);
nir_def_replace(&tex->def, result);
return true;
}

View file

@ -210,14 +210,7 @@ lower_builtin_instr(nir_builder *b, nir_intrinsic_instr *intrin,
def = nir_swizzle(b, def, swiz, intrin->num_components);
/* and rewrite uses of original instruction: */
nir_def_rewrite_uses(&intrin->def, def);
/* at this point intrin should be unused. We need to remove it
* (rather than waiting for DCE pass) to avoid dangling reference
* to remove'd var. And we have to remove the original uniform
* var since we don't want it to get uniform space allocated.
*/
nir_instr_remove(&intrin->instr);
nir_def_replace(&intrin->def, def);
return true;
}

View file

@ -237,8 +237,7 @@ clc_lower_input_image_deref(nir_builder *b, struct clc_image_lower_context *cont
}
/* No actual intrinsic needed here, just reference the loaded variable */
nir_def_rewrite_uses(&intrinsic->def, *cached_deref);
nir_instr_remove(&intrinsic->instr);
nir_def_replace(&intrinsic->def, *cached_deref);
break;
}

View file

@ -53,8 +53,7 @@ lower_load_base_global_invocation_id(nir_builder *b, nir_intrinsic_instr *intr,
nir_def *offset = load_ubo(b, intr, var, offsetof(struct clc_work_properties_data,
global_offset_x));
nir_def_rewrite_uses(&intr->def, offset);
nir_instr_remove(&intr->instr);
nir_def_replace(&intr->def, offset);
return true;
}
@ -66,8 +65,7 @@ lower_load_work_dim(nir_builder *b, nir_intrinsic_instr *intr,
nir_def *dim = load_ubo(b, intr, var, offsetof(struct clc_work_properties_data,
work_dim));
nir_def_rewrite_uses(&intr->def, dim);
nir_instr_remove(&intr->instr);
nir_def_replace(&intr->def, dim);
return true;
}
@ -80,8 +78,7 @@ lower_load_num_workgroups(nir_builder *b, nir_intrinsic_instr *intr,
nir_def *count =
load_ubo(b, intr, var, offsetof(struct clc_work_properties_data,
group_count_total_x));
nir_def_rewrite_uses(&intr->def, count);
nir_instr_remove(&intr->instr);
nir_def_replace(&intr->def, count);
return true;
}
@ -94,8 +91,7 @@ lower_load_base_workgroup_id(nir_builder *b, nir_intrinsic_instr *intr,
nir_def *offset =
load_ubo(b, intr, var, offsetof(struct clc_work_properties_data,
group_id_offset_x));
nir_def_rewrite_uses(&intr->def, offset);
nir_instr_remove(&intr->instr);
nir_def_replace(&intr->def, offset);
return true;
}
@ -177,8 +173,7 @@ lower_load_kernel_input(nir_builder *b, nir_intrinsic_instr *intr,
nir_def *result =
nir_load_deref(b, deref);
nir_def_rewrite_uses(&intr->def, result);
nir_instr_remove(&intr->instr);
nir_def_replace(&intr->def, result);
return true;
}

View file

@ -115,8 +115,7 @@ lower_32b_offset_load(nir_builder *b, nir_intrinsic_instr *intr, nir_variable *v
}
nir_def *result = nir_vec(b, comps, num_components);
nir_def_rewrite_uses(&intr->def, result);
nir_instr_remove(&intr->instr);
nir_def_replace(&intr->def, result);
return true;
}
@ -672,8 +671,7 @@ lower_shared_atomic(nir_builder *b, nir_intrinsic_instr *intr, nir_variable *var
result = nir_deref_atomic(b, 32, &deref->def, intr->src[1].ssa,
.atomic_op = nir_intrinsic_atomic_op(intr));
nir_def_rewrite_uses(&intr->def, result);
nir_instr_remove(&intr->instr);
nir_def_replace(&intr->def, result);
return true;
}
@ -757,9 +755,7 @@ lower_deref_ssbo(nir_builder *b, nir_deref_instr *deref)
nir_deref_instr *deref_cast =
nir_build_deref_cast(b, ptr, nir_var_mem_ssbo, deref->type,
glsl_get_explicit_stride(var->type));
nir_def_rewrite_uses(&deref->def,
&deref_cast->def);
nir_instr_remove(&deref->instr);
nir_def_replace(&deref->def, &deref_cast->def);
deref = deref_cast;
return true;
@ -884,8 +880,7 @@ cast_phi(nir_builder *b, nir_phi_instr *phi, unsigned new_bit_size)
b->cursor = nir_after_phis(nir_cursor_current_block(b->cursor));
nir_def *result = nir_u2uN(b, &lowered->def, old_bit_size);
nir_def_rewrite_uses(&phi->def, result);
nir_instr_remove(&phi->instr);
nir_def_replace(&phi->def, result);
}
static bool
@ -1225,8 +1220,7 @@ lower_load_local_group_size(nir_builder *b, nir_intrinsic_instr *intr)
nir_const_value_for_int(b->shader->info.workgroup_size[2], 32)
};
nir_def *size = nir_build_imm(b, 3, 32, v);
nir_def_rewrite_uses(&intr->def, size);
nir_instr_remove(&intr->instr);
nir_def_replace(&intr->def, size);
}
static bool
@ -2045,8 +2039,7 @@ split_unaligned_load(nir_builder *b, nir_intrinsic_instr *intrin, unsigned align
}
nir_def *new_dest = nir_extract_bits(b, srcs, num_loads, 0, num_comps, intrin->def.bit_size);
nir_def_rewrite_uses(&intrin->def, new_dest);
nir_instr_remove(&intrin->instr);
nir_def_replace(&intrin->def, new_dest);
}
static void
@ -2865,8 +2858,7 @@ kill_undefined_varyings(struct nir_builder *b,
*/
nir_def *zero = nir_imm_zero(b, intr->def.num_components,
intr->def.bit_size);
nir_def_rewrite_uses(&intr->def, zero);
nir_instr_remove(instr);
nir_def_replace(&intr->def, zero);
return true;
}

View file

@ -374,8 +374,7 @@ lower_load_push_constant(struct nir_builder *builder, nir_instr *instr,
.range_base = base,
.range = range);
nir_def_rewrite_uses(&intrin->def, load_data);
nir_instr_remove(instr);
nir_def_replace(&intrin->def, load_data);
return true;
}

View file

@ -80,8 +80,7 @@ nv50_nir_lower_load_user_clip_plane_cb(nir_builder *b, nir_intrinsic_instr *intr
nir_load_ubo(b, 4, 32, nir_imm_int(b, info->io.auxCBSlot),
nir_imm_int(b, offset), .range = ~0u);
nir_def_rewrite_uses(&intrin->def, replacement);
nir_instr_remove(&intrin->instr);
nir_def_replace(&intrin->def, replacement);
return true;
}

View file

@ -342,8 +342,7 @@ lower_fs_input_intrin(nir_builder *b, nir_intrinsic_instr *intrin, void *data)
return false;
}
nir_def_rewrite_uses(&intrin->def, res);
nir_instr_remove(&intrin->instr);
nir_def_replace(&intrin->def, res);
return true;
}

View file

@ -47,8 +47,7 @@ lower_ldcx_to_global(nir_builder *b, nir_intrinsic_instr *load)
nir_pop_if(b, NULL);
val = nir_if_phi(b, val, zero);
nir_def_rewrite_uses(&load->def, val);
nir_instr_remove(&load->instr);
nir_def_replace(&load->def, val);
}
struct non_uniform_section {

View file

@ -246,8 +246,7 @@ nak_nir_lower_scan_reduce_intrin(nir_builder *b,
data = nir_if_phi(b, full, partial);
}
nir_def_rewrite_uses(&intrin->def, data);
nir_instr_remove(&intrin->instr);
nir_def_replace(&intrin->def, data);
return true;
}

View file

@ -611,8 +611,7 @@ pan_inline_blend_constants(nir_builder *b, nir_intrinsic_instr *intr,
b->cursor = nir_after_instr(&intr->instr);
nir_def *constant = nir_build_imm(b, 4, 32, constants);
nir_def_rewrite_uses(&intr->def, constant);
nir_instr_remove(&intr->instr);
nir_def_replace(&intr->def, constant);
return true;
}

View file

@ -404,8 +404,7 @@ lower_res_intrinsic(nir_builder *b, nir_intrinsic_instr *intrin,
assert(intrin->def.bit_size == res->bit_size);
assert(intrin->def.num_components == res->num_components);
nir_def_rewrite_uses(&intrin->def, res);
nir_instr_remove(&intrin->instr);
nir_def_replace(&intrin->def, res);
return true;
}
@ -619,8 +618,7 @@ lower_tex(nir_builder *b, nir_tex_instr *tex, const struct lower_desc_ctx *ctx)
unreachable("Unsupported texture query op");
}
nir_def_rewrite_uses(&tex->def, res);
nir_instr_remove(&tex->instr);
nir_def_replace(&tex->def, res);
return true;
}
@ -718,8 +716,7 @@ lower_img_intrinsic(nir_builder *b, nir_intrinsic_instr *intr,
unreachable("Unsupported image query op");
}
nir_def_rewrite_uses(&intr->def, res);
nir_instr_remove(&intr->instr);
nir_def_replace(&intr->def, res);
} else {
nir_rewrite_image_intrinsic(intr, get_img_index(b, deref, ctx), false);
}

Some files were not shown because too many files have changed in this diff Show more