From c4b6b0d949e3f1210445bce1688dcd2395df3007 Mon Sep 17 00:00:00 2001 From: Erik Faye-Lund Date: Fri, 16 Jun 2023 15:38:39 +0200 Subject: [PATCH] intel: use imm-helpers Reviewed-by: Kenneth Graunke Part-of: --- src/gallium/drivers/crocus/crocus_program.c | 18 ++++++++-------- src/gallium/drivers/iris/iris_program.c | 14 ++++++------ src/intel/blorp/blorp_blit.c | 27 ++++++++++++------------ src/intel/blorp/blorp_clear.c | 3 +-- src/intel/compiler/brw_nir_lower_cs_intrinsics.c | 2 +- 5 files changed, 31 insertions(+), 33 deletions(-) diff --git a/src/gallium/drivers/crocus/crocus_program.c b/src/gallium/drivers/crocus/crocus_program.c index c3aa7ba..f1df95a 100644 --- a/src/gallium/drivers/crocus/crocus_program.c +++ b/src/gallium/drivers/crocus/crocus_program.c @@ -257,7 +257,7 @@ get_aoa_deref_offset(nir_builder *b, nir_ssa_def *index = nir_ssa_for_src(b, deref->arr.index, 1); assert(deref->arr.index.ssa); offset = nir_iadd(b, offset, - nir_imul(b, index, nir_imm_int(b, array_size))); + nir_imul_imm(b, index, array_size)); deref = nir_deref_instr_parent(deref); assert(glsl_type_is_array(deref->type)); @@ -301,8 +301,8 @@ crocus_lower_storage_image_derefs(nir_shader *nir) b.cursor = nir_before_instr(&intrin->instr); nir_ssa_def *index = - nir_iadd(&b, nir_imm_int(&b, var->data.driver_location), - get_aoa_deref_offset(&b, deref, 1)); + nir_iadd_imm(&b, get_aoa_deref_offset(&b, deref, 1), + var->data.driver_location); nir_rewrite_image_intrinsic(intrin, index, false); break; } @@ -597,10 +597,10 @@ crocus_setup_uniforms(ASSERTED const struct intel_device_info *devinfo, } b.cursor = nir_before_instr(instr); - offset = nir_iadd(&b, - get_aoa_deref_offset(&b, deref, BRW_IMAGE_PARAM_SIZE * 4), - nir_imm_int(&b, img_idx[var->data.binding] * 4 + - nir_intrinsic_base(intrin) * 16)); + offset = nir_iadd_imm(&b, + get_aoa_deref_offset(&b, deref, BRW_IMAGE_PARAM_SIZE * 4), + img_idx[var->data.binding] * 4 + + nir_intrinsic_base(intrin) * 16); break; } case nir_intrinsic_load_workgroup_size: { @@ -991,8 +991,8 @@ crocus_setup_binding_table(const struct intel_device_info *devinfo, nir_ssa_def *val = nir_fmul_imm(&b, &tex->dest.ssa, (1 << width) - 1); val = nir_f2u32(&b, val); if (wa & WA_SIGN) { - val = nir_ishl(&b, val, nir_imm_int(&b, 32 - width)); - val = nir_ishr(&b, val, nir_imm_int(&b, 32 - width)); + val = nir_ishl_imm(&b, val, 32 - width); + val = nir_ishr_imm(&b, val, 32 - width); } nir_ssa_def_rewrite_uses_after(&tex->dest.ssa, val, val->parent_instr); } diff --git a/src/gallium/drivers/iris/iris_program.c b/src/gallium/drivers/iris/iris_program.c index 95e225c..dcdf33c 100644 --- a/src/gallium/drivers/iris/iris_program.c +++ b/src/gallium/drivers/iris/iris_program.c @@ -240,7 +240,7 @@ get_aoa_deref_offset(nir_builder *b, nir_ssa_def *index = nir_ssa_for_src(b, deref->arr.index, 1); assert(deref->arr.index.ssa); offset = nir_iadd(b, offset, - nir_imul(b, index, nir_imm_int(b, array_size))); + nir_imul_imm(b, index, array_size)); deref = nir_deref_instr_parent(deref); assert(glsl_type_is_array(deref->type)); @@ -284,8 +284,8 @@ iris_lower_storage_image_derefs(nir_shader *nir) b.cursor = nir_before_instr(&intrin->instr); nir_ssa_def *index = - nir_iadd(&b, nir_imm_int(&b, var->data.driver_location), - get_aoa_deref_offset(&b, deref, 1)); + nir_iadd_imm(&b, get_aoa_deref_offset(&b, deref, 1), + var->data.driver_location); nir_rewrite_image_intrinsic(intrin, index, false); break; } @@ -624,11 +624,11 @@ iris_setup_uniforms(ASSERTED const struct intel_device_info *devinfo, } b.cursor = nir_before_instr(instr); - offset = nir_iadd(&b, + offset = nir_iadd_imm(&b, get_aoa_deref_offset(&b, deref, BRW_IMAGE_PARAM_SIZE * 4), - nir_imm_int(&b, system_values_start + - img_idx[var->data.binding] * 4 + - nir_intrinsic_base(intrin) * 16)); + system_values_start + + img_idx[var->data.binding] * 4 + + nir_intrinsic_base(intrin) * 16); break; } case nir_intrinsic_load_workgroup_size: { diff --git a/src/intel/blorp/blorp_blit.c b/src/intel/blorp/blorp_blit.c index d948e1c..62b3d3c 100644 --- a/src/intel/blorp/blorp_blit.c +++ b/src/intel/blorp/blorp_blit.c @@ -818,22 +818,21 @@ blorp_nir_manual_blend_bilinear(nir_builder *b, nir_ssa_def *pos, sample = nir_f2i32(b, sample); if (tex_samples == 2) { - sample = nir_isub(b, nir_imm_int(b, 1), sample); + sample = nir_isub_imm(b, 1, sample); } else if (tex_samples == 8) { - sample = nir_iand(b, nir_ishr(b, nir_imm_int(b, 0x64210573), - nir_ishl(b, sample, nir_imm_int(b, 2))), - nir_imm_int(b, 0xf)); + sample = nir_iand_imm(b, nir_ishr(b, nir_imm_int(b, 0x64210573), + nir_ishl_imm(b, sample, 2)), + 0xf); } else if (tex_samples == 16) { nir_ssa_def *sample_low = - nir_iand(b, nir_ishr(b, nir_imm_int(b, 0xd31479af), - nir_ishl(b, sample, nir_imm_int(b, 2))), - nir_imm_int(b, 0xf)); + nir_iand_imm(b, nir_ishr(b, nir_imm_int(b, 0xd31479af), + nir_ishl_imm(b, sample, 2)), + 0xf); nir_ssa_def *sample_high = - nir_iand(b, nir_ishr(b, nir_imm_int(b, 0xe58b602c), - nir_ishl(b, nir_iadd(b, sample, - nir_imm_int(b, -8)), - nir_imm_int(b, 2))), - nir_imm_int(b, 0xf)); + nir_iand_imm(b, nir_ishr(b, nir_imm_int(b, 0xe58b602c), + nir_ishl_imm(b, nir_iadd_imm(b, sample, -8), + 2)), + 0xf); sample = nir_bcsel(b, nir_ilt_imm(b, sample, 8), sample_low, sample_high); @@ -905,8 +904,8 @@ bit_cast_color(struct nir_builder *b, nir_ssa_def *color, const unsigned chan_start_bit = dst_fmtl->channels_array[c].start_bit; const unsigned chan_bits = dst_fmtl->channels_array[c].bits; - chans[c] = nir_iand(b, nir_shift_imm(b, packed, -(int)chan_start_bit), - nir_imm_int(b, BITFIELD_MASK(chan_bits))); + chans[c] = nir_iand_imm(b, nir_shift_imm(b, packed, -(int)chan_start_bit), + BITFIELD_MASK(chan_bits)); if (dst_fmtl->channels_array[c].type == ISL_UNORM) chans[c] = nir_format_unorm_to_float(b, chans[c], &chan_bits); diff --git a/src/intel/blorp/blorp_clear.c b/src/intel/blorp/blorp_clear.c index 89ae7f8..c3a6f9c 100644 --- a/src/intel/blorp/blorp_clear.c +++ b/src/intel/blorp/blorp_clear.c @@ -1311,8 +1311,7 @@ blorp_ccs_resolve(struct blorp_batch *batch, static nir_ssa_def * blorp_nir_bit(nir_builder *b, nir_ssa_def *src, unsigned bit) { - return nir_iand(b, nir_ushr(b, src, nir_imm_int(b, bit)), - nir_imm_int(b, 1)); + return nir_iand_imm(b, nir_ushr_imm(b, src, bit), 1); } #pragma pack(push, 1) diff --git a/src/intel/compiler/brw_nir_lower_cs_intrinsics.c b/src/intel/compiler/brw_nir_lower_cs_intrinsics.c index f5a7b58..679c44b 100644 --- a/src/intel/compiler/brw_nir_lower_cs_intrinsics.c +++ b/src/intel/compiler/brw_nir_lower_cs_intrinsics.c @@ -100,7 +100,7 @@ compute_local_index_id(nir_builder *b, id_x = nir_umod(b, block, size_x); id_y = nir_umod(b, nir_iadd(b, - nir_umod(b, linear, nir_imm_int(b, height)), + nir_umod_imm(b, linear, height), nir_imul_imm(b, nir_udiv(b, block, size_x), height)), -- 2.7.4