OpenHarmony/third_party_mesa3d
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

intel: use imm-helpers

Browse Source 
Reviewed-by: Kenneth Graunke <kenneth@whitecape.org>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/23855>
This commit is contained in:
Erik Faye-Lund
2023-06-16 15:38:39 +02:00
committed by Marge Bot
parent 0b57f76986
commit c4b6b0d949
5 changed files with 31 additions and 33 deletions
Download Patch File Download Diff File Expand all files Collapse all files

18
src/gallium/drivers/crocus/crocus_program.c
View File

@@ -257,7 +257,7 @@ get_aoa_deref_offset(nir_builder *b,
nir_ssa_def *index = nir_ssa_for_src(b, deref->arr.index, 1); nir_ssa_def *index = nir_ssa_for_src(b, deref->arr.index, 1);
assert(deref->arr.index.ssa); assert(deref->arr.index.ssa);
offset = nir_iadd(b, offset, 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); deref = nir_deref_instr_parent(deref);
assert(glsl_type_is_array(deref->type)); 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); b.cursor = nir_before_instr(&intrin->instr);
nir_ssa_def *index = nir_ssa_def *index =
nir_iadd(&b, nir_imm_int(&b, var->data.driver_location), nir_iadd_imm(&b, get_aoa_deref_offset(&b, deref, 1),
get_aoa_deref_offset(&b, deref, 1)); var->data.driver_location);
nir_rewrite_image_intrinsic(intrin, index, false); nir_rewrite_image_intrinsic(intrin, index, false);
break; break;
} }
@@ -597,10 +597,10 @@ crocus_setup_uniforms(ASSERTED const struct intel_device_info *devinfo,
} }
b.cursor = nir_before_instr(instr); 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), get_aoa_deref_offset(&b, deref, BRW_IMAGE_PARAM_SIZE * 4),
nir_imm_int(&b, img_idx[var->data.binding] * 4 + img_idx[var->data.binding] * 4 +
nir_intrinsic_base(intrin) * 16)); nir_intrinsic_base(intrin) * 16);
break; break;
} }
case nir_intrinsic_load_workgroup_size: { 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); nir_ssa_def *val = nir_fmul_imm(&b, &tex->dest.ssa, (1 << width) - 1);
val = nir_f2u32(&b, val); val = nir_f2u32(&b, val);
if (wa & WA_SIGN) { if (wa & WA_SIGN) {
val = nir_ishl(&b, val, nir_imm_int(&b, 32 - width)); val = nir_ishl_imm(&b, val, 32 - width);
val = nir_ishr(&b, val, nir_imm_int(&b, 32 - width)); val = nir_ishr_imm(&b, val, 32 - width);
} }
nir_ssa_def_rewrite_uses_after(&tex->dest.ssa, val, val->parent_instr); nir_ssa_def_rewrite_uses_after(&tex->dest.ssa, val, val->parent_instr);
} }

14
src/gallium/drivers/iris/iris_program.c
View File

@@ -240,7 +240,7 @@ get_aoa_deref_offset(nir_builder *b,
nir_ssa_def *index = nir_ssa_for_src(b, deref->arr.index, 1); nir_ssa_def *index = nir_ssa_for_src(b, deref->arr.index, 1);
assert(deref->arr.index.ssa); assert(deref->arr.index.ssa);
offset = nir_iadd(b, offset, 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); deref = nir_deref_instr_parent(deref);
assert(glsl_type_is_array(deref->type)); 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); b.cursor = nir_before_instr(&intrin->instr);
nir_ssa_def *index = nir_ssa_def *index =
nir_iadd(&b, nir_imm_int(&b, var->data.driver_location), nir_iadd_imm(&b, get_aoa_deref_offset(&b, deref, 1),
get_aoa_deref_offset(&b, deref, 1)); var->data.driver_location);
nir_rewrite_image_intrinsic(intrin, index, false); nir_rewrite_image_intrinsic(intrin, index, false);
break; break;
} }
@@ -624,11 +624,11 @@ iris_setup_uniforms(ASSERTED const struct intel_device_info *devinfo,
} }
b.cursor = nir_before_instr(instr); 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), get_aoa_deref_offset(&b, deref, BRW_IMAGE_PARAM_SIZE * 4),
nir_imm_int(&b, system_values_start + system_values_start +
img_idx[var->data.binding] * 4 + img_idx[var->data.binding] * 4 +
nir_intrinsic_base(intrin) * 16)); nir_intrinsic_base(intrin) * 16);
break; break;
} }
case nir_intrinsic_load_workgroup_size: { case nir_intrinsic_load_workgroup_size: {

27
src/intel/blorp/blorp_blit.c
View File

@@ -818,22 +818,21 @@ blorp_nir_manual_blend_bilinear(nir_builder *b, nir_ssa_def *pos,
sample = nir_f2i32(b, sample); sample = nir_f2i32(b, sample);
if (tex_samples == 2) { 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) { } else if (tex_samples == 8) {
sample = nir_iand(b, nir_ishr(b, nir_imm_int(b, 0x64210573), sample = nir_iand_imm(b, nir_ishr(b, nir_imm_int(b, 0x64210573),
nir_ishl(b, sample, nir_imm_int(b, 2))), nir_ishl_imm(b, sample, 2)),
nir_imm_int(b, 0xf)); 0xf);
} else if (tex_samples == 16) { } else if (tex_samples == 16) {
nir_ssa_def *sample_low = nir_ssa_def *sample_low =
nir_iand(b, nir_ishr(b, nir_imm_int(b, 0xd31479af), nir_iand_imm(b, nir_ishr(b, nir_imm_int(b, 0xd31479af),
nir_ishl(b, sample, nir_imm_int(b, 2))), nir_ishl_imm(b, sample, 2)),
nir_imm_int(b, 0xf)); 0xf);
nir_ssa_def *sample_high = nir_ssa_def *sample_high =
nir_iand(b, nir_ishr(b, nir_imm_int(b, 0xe58b602c), nir_iand_imm(b, nir_ishr(b, nir_imm_int(b, 0xe58b602c),
nir_ishl(b, nir_iadd(b, sample, nir_ishl_imm(b, nir_iadd_imm(b, sample, -8),
nir_imm_int(b, -8)), 2)),
nir_imm_int(b, 2))), 0xf);
nir_imm_int(b, 0xf));
sample = nir_bcsel(b, nir_ilt_imm(b, sample, 8), sample = nir_bcsel(b, nir_ilt_imm(b, sample, 8),
sample_low, sample_high); 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_start_bit = dst_fmtl->channels_array[c].start_bit;
const unsigned chan_bits = dst_fmtl->channels_array[c].bits; const unsigned chan_bits = dst_fmtl->channels_array[c].bits;
chans[c] = nir_iand(b, nir_shift_imm(b, packed, -(int)chan_start_bit), chans[c] = nir_iand_imm(b, nir_shift_imm(b, packed, -(int)chan_start_bit),
nir_imm_int(b, BITFIELD_MASK(chan_bits))); BITFIELD_MASK(chan_bits));
if (dst_fmtl->channels_array[c].type == ISL_UNORM) if (dst_fmtl->channels_array[c].type == ISL_UNORM)
chans[c] = nir_format_unorm_to_float(b, chans[c], &chan_bits); chans[c] = nir_format_unorm_to_float(b, chans[c], &chan_bits);

3
src/intel/blorp/blorp_clear.c
View File

@@ -1311,8 +1311,7 @@ blorp_ccs_resolve(struct blorp_batch *batch,
static nir_ssa_def * static nir_ssa_def *
blorp_nir_bit(nir_builder *b, nir_ssa_def *src, unsigned bit) 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)), return nir_iand_imm(b, nir_ushr_imm(b, src, bit), 1);
nir_imm_int(b, 1));
} }
#pragma pack(push, 1) #pragma pack(push, 1)

2
src/intel/compiler/brw_nir_lower_cs_intrinsics.c
View File

@@ -100,7 +100,7 @@ compute_local_index_id(nir_builder *b,
id_x = nir_umod(b, block, size_x); id_x = nir_umod(b, block, size_x);
id_y = nir_umod(b, id_y = nir_umod(b,
nir_iadd(b, nir_iadd(b,
nir_umod(b, linear, nir_imm_int(b, height)), nir_umod_imm(b, linear, height),
nir_imul_imm(b, nir_imul_imm(b,
nir_udiv(b, block, size_x), nir_udiv(b, block, size_x),
height)), height)),