intel: use imm-helpers

Reviewed-by: Kenneth Graunke <kenneth@whitecape.org>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/23855>

diff --git a/src/gallium/drivers/crocus/crocus_program.c b/src/gallium/drivers/crocus/crocus_program.c
index c3aa7ba..f1df95a 100644 (file)
--- 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 (file)
--- 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 (file)
--- 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 (file)
--- 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 (file)
--- 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)),