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

nir/builder: Drop nir_i2i and nir_u2u in favor of nir_x2xN

Browse Source 
Reviewed-by: Alyssa Rosenzweig <alyssa.rosenzweig@collabora.com>
Reviewed-by: Emma Anholt <emma@anholt.net>
Reviewed-by: Timur Kristóf <timur.kristof@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/20067>
This commit is contained in:
Jason Ekstrand
2022-11-29 12:54:21 -06:00
committed by Marge Bot
parent ccf19e0956
commit d9a24632d3
19 changed files with 43 additions and 73 deletions
Download Patch File Download Diff File Expand all files Collapse all files

6
src/amd/common/ac_nir_lower_ngg.c
View File

@@ -2794,7 +2794,7 @@ ngg_gs_export_vertices(nir_builder *b, nir_ssa_def *max_num_out_vtx, nir_ssa_def
/* Convert to the expected bit size of the output variable. */
unsigned bit_size = glsl_base_type_bit_size(glsl_get_base_type(var->type));
if (bit_size != 32)
val = nir_u2u(b, val, bit_size);
val = nir_u2uN(b, val, bit_size);
nir_store_output(b, val, nir_imm_int(b, 0), .base = info->base,
.io_semantics = io_sem, .component = start + i, .write_mask = 1);
@@ -3397,7 +3397,7 @@ lower_ms_load_output(nir_builder *b,
} else if (io_sem.location == VARYING_SLOT_PRIMITIVE_INDICES) {
nir_ssa_def *offset_src = nir_get_io_offset_src(intrin)->ssa;
nir_ssa_def *index = ms_load_prim_indices(b, offset_src, s);
return nir_u2u(b, index, intrin->dest.ssa.bit_size);
return nir_u2uN(b, index, intrin->dest.ssa.bit_size);
}
unreachable("Invalid mesh shader output");
@@ -3495,7 +3495,7 @@ regroup_load_val(nir_builder *b, nir_ssa_def *load, unsigned dest_bit_size)
assert(num_components <= 4);
nir_ssa_def *components[4] = {0};
for (unsigned i = 0; i < num_components; ++i)
components[i] = nir_u2u(b, nir_channel(b, load, i), dest_bit_size);
components[i] = nir_u2uN(b, nir_channel(b, load, i), dest_bit_size);
return nir_vec(b, components, num_components);
}

36
src/compiler/nir/nir_builder.h
View File

@@ -669,36 +669,6 @@ nir_vector_insert(nir_builder *b, nir_ssa_def *vec, nir_ssa_def *scalar,
}
}
static inline nir_ssa_def *
nir_i2i(nir_builder *build, nir_ssa_def *x, unsigned dest_bit_size)
{
if (x->bit_size == dest_bit_size)
return x;
switch (dest_bit_size) {
case 64: return nir_i2i64(build, x);
case 32: return nir_i2i32(build, x);
case 16: return nir_i2i16(build, x);
case 8: return nir_i2i8(build, x);
default: unreachable("Invalid bit size");
}
}
static inline nir_ssa_def *
nir_u2u(nir_builder *build, nir_ssa_def *x, unsigned dest_bit_size)
{
if (x->bit_size == dest_bit_size)
return x;
switch (dest_bit_size) {
case 64: return nir_u2u64(build, x);
case 32: return nir_u2u32(build, x);
case 16: return nir_u2u16(build, x);
case 8: return nir_u2u8(build, x);
default: unreachable("Invalid bit size");
}
}
static inline nir_ssa_def *
nir_iadd_imm(nir_builder *build, nir_ssa_def *x, uint64_t y)
{
@@ -963,7 +933,7 @@ nir_pack_bits(nir_builder *b, nir_ssa_def *src, unsigned dest_bit_size)
/* If we got here, we have no dedicated unpack opcode. */
nir_ssa_def *dest = nir_imm_intN_t(b, 0, dest_bit_size);
for (unsigned i = 0; i < src->num_components; i++) {
nir_ssa_def *val = nir_u2u(b, nir_channel(b, src, i), dest_bit_size);
nir_ssa_def *val = nir_u2uN(b, nir_channel(b, src, i), dest_bit_size);
val = nir_ishl(b, val, nir_imm_int(b, i * src->bit_size));
dest = nir_ior(b, dest, val);
}
@@ -1000,7 +970,7 @@ nir_unpack_bits(nir_builder *b, nir_ssa_def *src, unsigned dest_bit_size)
nir_ssa_def *dest_comps[NIR_MAX_VEC_COMPONENTS];
for (unsigned i = 0; i < dest_num_components; i++) {
nir_ssa_def *val = nir_ushr_imm(b, src, i * dest_bit_size);
dest_comps[i] = nir_u2u(b, val, dest_bit_size);
dest_comps[i] = nir_u2uN(b, val, dest_bit_size);
}
return nir_vec(b, dest_comps, dest_num_components);
}
@@ -1395,7 +1365,7 @@ nir_build_deref_follower(nir_builder *b, nir_deref_instr *parent,
if (leader->deref_type == nir_deref_type_array) {
assert(leader->arr.index.is_ssa);
nir_ssa_def *index = nir_i2i(b, leader->arr.index.ssa,
nir_ssa_def *index = nir_i2iN(b, leader->arr.index.ssa,
parent->dest.ssa.bit_size);
return nir_build_deref_array(b, parent, index);
} else {

4
src/compiler/nir/nir_builtin_builder.h
View File

@@ -238,7 +238,7 @@ nir_clz_u(nir_builder *b, nir_ssa_def *a)
{
nir_ssa_def *val;
val = nir_isub(b, nir_imm_intN_t(b, a->bit_size - 1, 32), nir_ufind_msb(b, a));
return nir_u2u(b, val, a->bit_size);
return nir_u2uN(b, val, a->bit_size);
}
static inline nir_ssa_def *
@@ -248,7 +248,7 @@ nir_ctz_u(nir_builder *b, nir_ssa_def *a)
return nir_bcsel(b, cond,
nir_imm_intN_t(b, a->bit_size, a->bit_size),
nir_u2u(b, nir_find_lsb(b, a), a->bit_size));
nir_u2uN(b, nir_find_lsb(b, a), a->bit_size));
}
#ifdef __cplusplus

6
src/compiler/nir/nir_lower_bit_size.c
View File

@@ -199,7 +199,7 @@ lower_intrinsic_instr(nir_builder *b, nir_intrinsic_instr *intrin,
if (intrin->intrinsic != nir_intrinsic_vote_feq &&
intrin->intrinsic != nir_intrinsic_vote_ieq)
res = nir_u2u(b, res, old_bit_size);
res = nir_u2uN(b, res, old_bit_size);
nir_ssa_def_rewrite_uses(&intrin->dest.ssa, res);
break;
@@ -221,7 +221,7 @@ lower_phi_instr(nir_builder *b, nir_phi_instr *phi, unsigned bit_size,
nir_foreach_phi_src(src, phi) {
b->cursor = nir_after_block_before_jump(src->pred);
assert(src->src.is_ssa);
nir_ssa_def *new_src = nir_u2u(b, src->src.ssa, bit_size);
nir_ssa_def *new_src = nir_u2uN(b, src->src.ssa, bit_size);
nir_instr_rewrite_src(&phi->instr, &src->src, nir_src_for_ssa(new_src));
}
@@ -230,7 +230,7 @@ lower_phi_instr(nir_builder *b, nir_phi_instr *phi, unsigned bit_size,
b->cursor = nir_after_instr(&last_phi->instr);
nir_ssa_def *new_dest = nir_u2u(b, &phi->dest.ssa, old_bit_size);
nir_ssa_def *new_dest = nir_u2uN(b, &phi->dest.ssa, old_bit_size);
nir_ssa_def_rewrite_uses_after(&phi->dest.ssa, new_dest,
new_dest->parent_instr);
}

4
src/compiler/nir/nir_lower_io.c
View File

@@ -1882,9 +1882,9 @@ nir_explicit_io_address_from_deref(nir_builder *b, nir_deref_instr *deref,
*/
if (deref->arr.in_bounds && deref->deref_type == nir_deref_type_array) {
index = nir_u2u32(b, index);
offset = nir_u2u(b, nir_amul_imm(b, index, stride), offset_bit_size);
offset = nir_u2uN(b, nir_amul_imm(b, index, stride), offset_bit_size);
} else {
index = nir_i2i(b, index, offset_bit_size);
index = nir_i2iN(b, index, offset_bit_size);
offset = nir_amul_imm(b, index, stride);
}

2
src/compiler/nir/nir_lower_io_to_vector.c
View File

@@ -343,7 +343,7 @@ build_array_index(nir_builder *b, nir_deref_instr *deref, nir_ssa_def *base,
case nir_deref_type_var:
return base;
case nir_deref_type_array: {
nir_ssa_def *index = nir_i2i(b, deref->arr.index.ssa,
nir_ssa_def *index = nir_i2iN(b, deref->arr.index.ssa,
deref->dest.ssa.bit_size);
if (nir_deref_instr_parent(deref)->deref_type == nir_deref_type_var &&

2
src/compiler/nir/nir_lower_locals_to_regs.c
View File

@@ -166,7 +166,7 @@ get_deref_reg_src(nir_deref_instr *deref, struct locals_to_regs_state *state)
}
assert(src.reg.indirect->is_ssa);
nir_ssa_def *index = nir_i2i(b, nir_ssa_for_src(b, d->arr.index, 1), 32);
nir_ssa_def *index = nir_i2iN(b, nir_ssa_for_src(b, d->arr.index, 1), 32);
src.reg.indirect->ssa =
nir_iadd(b, src.reg.indirect->ssa,
nir_imul_imm(b, index, inner_array_size));

4
src/compiler/nir/nir_lower_memcpy.c
View File

@@ -50,7 +50,7 @@ memcpy_load_deref_elem(nir_builder *b, nir_deref_instr *parent,
{
nir_deref_instr *deref;
index = nir_i2i(b, index, nir_dest_bit_size(parent->dest));
index = nir_i2iN(b, index, nir_dest_bit_size(parent->dest));
assert(parent->deref_type == nir_deref_type_cast);
deref = nir_build_deref_ptr_as_array(b, parent, index);
@@ -71,7 +71,7 @@ memcpy_store_deref_elem(nir_builder *b, nir_deref_instr *parent,
{
nir_deref_instr *deref;
index = nir_i2i(b, index, nir_dest_bit_size(parent->dest));
index = nir_i2iN(b, index, nir_dest_bit_size(parent->dest));
assert(parent->deref_type == nir_deref_type_cast);
deref = nir_build_deref_ptr_as_array(b, parent, index);
nir_store_deref(b, deref, value, ~0);

4
src/compiler/nir/nir_lower_printf.c
View File

@@ -82,7 +82,7 @@ lower_printf_instr(nir_builder *b, nir_instr *instr, void *_options)
/* Write the format string ID */
nir_ssa_def *fmt_str_id_offset =
nir_i2i(b, offset, ptr_bit_size);
nir_i2iN(b, offset, ptr_bit_size);
nir_deref_instr *fmt_str_id_deref =
nir_build_deref_array(b, buffer, fmt_str_id_offset);
fmt_str_id_deref = nir_build_deref_cast(b, &fmt_str_id_deref->dest.ssa,
@@ -111,7 +111,7 @@ lower_printf_instr(nir_builder *b, nir_instr *instr, void *_options)
unsigned field_offset = glsl_get_struct_field_offset(args->type, i);
nir_ssa_def *arg_offset =
nir_i2i(b, nir_iadd_imm(b, offset,
nir_i2iN(b, nir_iadd_imm(b, offset,
fmt_str_id_size + field_offset),
ptr_bit_size);
nir_deref_instr *dst_arg_deref =

18
src/compiler/nir/nir_lower_system_values.c
View File

@@ -48,7 +48,7 @@ sanitize_32bit_sysval(nir_builder *b, nir_intrinsic_instr *intrin)
return NULL;
intrin->dest.ssa.bit_size = 32;
return nir_u2u(b, &intrin->dest.ssa, bit_size);
return nir_u2uN(b, &intrin->dest.ssa, bit_size);
}
static nir_ssa_def*
@@ -56,7 +56,7 @@ build_global_group_size(nir_builder *b, unsigned bit_size)
{
nir_ssa_def *group_size = nir_load_workgroup_size(b);
nir_ssa_def *num_workgroups = nir_load_num_workgroups(b, bit_size);
return nir_imul(b, nir_u2u(b, group_size, bit_size),
return nir_imul(b, nir_u2uN(b, group_size, bit_size),
num_workgroups);
}
@@ -290,7 +290,7 @@ lower_id_to_index_no_umod(nir_builder *b, nir_ssa_def *index,
nir_ssa_def *y_portion = nir_imul(b, id_y, size_x);
nir_ssa_def *id_x = nir_isub(b, index, nir_iadd(b, z_portion, y_portion));
return nir_u2u(b, nir_vec3(b, id_x, id_y, id_z), bit_size);
return nir_u2uN(b, nir_vec3(b, id_x, id_y, id_z), bit_size);
}
@@ -321,7 +321,7 @@ lower_id_to_index(nir_builder *b, nir_ssa_def *index, nir_ssa_def *size,
nir_ssa_def *id_y = nir_umod(b, nir_udiv(b, index, size_x), size_y);
nir_ssa_def *id_z = nir_udiv(b, index, nir_imul(b, size_x, size_y));
return nir_u2u(b, nir_vec3(b, id_x, id_y, id_z), bit_size);
return nir_u2uN(b, nir_vec3(b, id_x, id_y, id_z), bit_size);
}
static bool
@@ -501,7 +501,7 @@ lower_compute_system_value_instr(nir_builder *b,
index = nir_iadd(b, index,
nir_imul(b, nir_channel(b, local_id, 1), size_x));
index = nir_iadd(b, index, nir_channel(b, local_id, 0));
return nir_u2u(b, index, bit_size);
return nir_u2uN(b, index, bit_size);
} else {
return NULL;
}
@@ -521,7 +521,7 @@ lower_compute_system_value_instr(nir_builder *b,
workgroup_size_const[0].u32 = b->shader->info.workgroup_size[0];
workgroup_size_const[1].u32 = b->shader->info.workgroup_size[1];
workgroup_size_const[2].u32 = b->shader->info.workgroup_size[2];
return nir_u2u(b, nir_build_imm(b, 3, 32, workgroup_size_const), bit_size);
return nir_u2uN(b, nir_build_imm(b, 3, 32, workgroup_size_const), bit_size);
}
case nir_intrinsic_load_global_invocation_id_zero_base: {
@@ -532,8 +532,8 @@ lower_compute_system_value_instr(nir_builder *b,
nir_ssa_def *local_id = nir_load_local_invocation_id(b);
return nir_iadd(b, nir_imul(b, group_id,
nir_u2u(b, group_size, bit_size)),
nir_u2u(b, local_id, bit_size));
nir_u2uN(b, group_size, bit_size)),
nir_u2uN(b, local_id, bit_size));
} else {
return NULL;
}
@@ -569,7 +569,7 @@ lower_compute_system_value_instr(nir_builder *b,
case nir_intrinsic_load_workgroup_id: {
if (options && options->has_base_workgroup_id)
return nir_iadd(b, nir_u2u(b, nir_load_workgroup_id_zero_base(b), bit_size),
return nir_iadd(b, nir_u2uN(b, nir_load_workgroup_id_zero_base(b), bit_size),
nir_load_base_workgroup_id(b, bit_size));
else if (options && options->lower_workgroup_id_to_index)
return lower_id_to_index_no_umod(b, nir_load_workgroup_index(b),

4
src/compiler/spirv/spirv_to_nir.c
View File

@@ -3593,7 +3593,7 @@ vtn_handle_image(struct vtn_builder *b, SpvOp opcode,
if (opcode == SpvOpImageQuerySize ||
opcode == SpvOpImageQuerySizeLod)
result = nir_u2u(&b->nb, result, glsl_get_bit_size(type->type));
result = nir_u2uN(&b->nb, result, glsl_get_bit_size(type->type));
if (opcode == SpvOpImageSparseRead) {
struct vtn_ssa_value *dest = vtn_create_ssa_value(b, struct_type->type);
@@ -5540,7 +5540,7 @@ vtn_handle_ptr(struct vtn_builder *b, SpvOp opcode,
vtn_get_nir_ssa(b, w[4]),
addr_format);
def = nir_idiv(&b->nb, def, nir_imm_intN_t(&b->nb, elem_size, def->bit_size));
def = nir_i2i(&b->nb, def, glsl_get_bit_size(type));
def = nir_i2iN(&b->nb, def, glsl_get_bit_size(type));
break;
}

10
src/compiler/spirv/vtn_alu.c
View File

@@ -929,7 +929,7 @@ vtn_handle_alu(struct vtn_builder *b, SpvOp opcode,
/* bit_count always returns int32, but the SPIR-V opcode just says the return
* value needs to be big enough to store the number of bits.
*/
dest->def = nir_u2u(&b->nb, nir_bit_count(&b->nb, src[0]), glsl_get_bit_size(dest_type));
dest->def = nir_u2uN(&b->nb, nir_bit_count(&b->nb, src[0]), glsl_get_bit_size(dest_type));
break;
}
@@ -1270,12 +1270,12 @@ vtn_handle_integer_dot(struct vtn_builder *b, SpvOp opcode,
*/
if (num_inputs == 3) {
dest = is_signed
? nir_iadd_sat(&b->nb, nir_i2i(&b->nb, dest, dest_size), src[2])
: nir_uadd_sat(&b->nb, nir_u2u(&b->nb, dest, dest_size), src[2]);
? nir_iadd_sat(&b->nb, nir_i2iN(&b->nb, dest, dest_size), src[2])
: nir_uadd_sat(&b->nb, nir_u2uN(&b->nb, dest, dest_size), src[2]);
} else {
dest = is_signed
? nir_i2i(&b->nb, dest, dest_size)
: nir_u2u(&b->nb, dest, dest_size);
? nir_i2iN(&b->nb, dest, dest_size)
: nir_u2uN(&b->nb, dest, dest_size);
}
}
}

2
src/compiler/spirv/vtn_opencl.c
View File

@@ -279,7 +279,7 @@ handle_alu(struct vtn_builder *b, uint32_t opcode,
nir_ssa_def *ret = nir_build_alu(&b->nb, nir_alu_op_for_opencl_opcode(b, (enum OpenCLstd_Entrypoints)opcode),
srcs[0], srcs[1], srcs[2], NULL);
if (opcode == OpenCLstd_Popcount)
ret = nir_u2u(&b->nb, ret, glsl_get_bit_size(dest_type->type));
ret = nir_u2uN(&b->nb, ret, glsl_get_bit_size(dest_type->type));
return ret;
}

2
src/compiler/spirv/vtn_variables.c
View File

@@ -206,7 +206,7 @@ vtn_access_link_as_ssa(struct vtn_builder *b, struct vtn_access_link link,
} else {
nir_ssa_def *ssa = vtn_ssa_value(b, link.id)->def;
if (ssa->bit_size != bit_size)
ssa = nir_i2i(&b->nb, ssa, bit_size);
ssa = nir_i2iN(&b->nb, ssa, bit_size);
return nir_imul_imm(&b->nb, ssa, stride);
}
}

2
src/gallium/frontends/clover/nir/invocation.cpp
View File

@@ -162,7 +162,7 @@ clover_lower_nir_instr(nir_builder *b, nir_instr *instr, void *_state)
loads[i] = var ? nir_load_var(b, var) : nir_imm_int(b, 0);
}
return nir_u2u(b, nir_vec(b, loads, state->global_dims),
return nir_u2uN(b, nir_vec(b, loads, state->global_dims),
nir_dest_bit_size(intrinsic->dest));
}
case nir_intrinsic_load_constant_base_ptr: {

4
src/gallium/frontends/rusticl/rusticl_nir.c
View File

@@ -54,7 +54,7 @@ rusticl_lower_intrinsics_instr(
deref = nir_build_deref_var(b, var);
deref = nir_build_deref_array(b, deref, val);
val = nir_u2u(b, nir_load_deref(b, deref), 32);
val = nir_u2uN(b, nir_load_deref(b, deref), 32);
// we have to fix up the value base
val = nir_iadd_imm(b, val, -offset);
@@ -69,7 +69,7 @@ rusticl_lower_intrinsics_instr(
return nir_load_var(b, state->printf_buf);
case nir_intrinsic_load_work_dim:
assert(state->work_dim);
return nir_u2u(b, nir_load_var(b, state->work_dim), nir_dest_bit_size(intrins->dest));
return nir_u2uN(b, nir_load_var(b, state->work_dim), nir_dest_bit_size(intrins->dest));
default:
return NULL;
}

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

@@ -226,7 +226,7 @@ lower_kernel_intrinsics(nir_shader *nir)
/* We may need to do a bit-size cast here */
nir_ssa_def *num_work_groups =
nir_u2u(&b, &load->dest.ssa, intrin->dest.ssa.bit_size);
nir_u2uN(&b, &load->dest.ssa, intrin->dest.ssa.bit_size);
nir_ssa_def_rewrite_uses(&intrin->dest.ssa, num_work_groups);
progress = true;

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

@@ -43,7 +43,7 @@ resize_deref(nir_builder *b, nir_deref_instr *deref,
if (nir_src_is_const(deref->arr.index)) {
idx = nir_imm_intN_t(b, nir_src_as_int(deref->arr.index), bit_size);
} else {
idx = nir_i2i(b, deref->arr.index.ssa, bit_size);
idx = nir_i2iN(b, deref->arr.index.ssa, bit_size);
}
nir_instr_rewrite_src(&deref->instr, &deref->arr.index,
nir_src_for_ssa(idx));

2
src/microsoft/clc/clc_nir.c
View File

@@ -172,7 +172,7 @@ lower_load_kernel_input(nir_builder *b, nir_intrinsic_instr *intr,
const struct glsl_type *type =
glsl_vector_type(base_type, nir_dest_num_components(intr->dest));
nir_ssa_def *ptr = nir_vec2(b, nir_imm_int(b, var->data.binding),
nir_u2u(b, intr->src[0].ssa, 32));
nir_u2uN(b, intr->src[0].ssa, 32));
nir_deref_instr *deref = nir_build_deref_cast(b, ptr, nir_var_mem_ubo, type,
bit_size / 8);
deref->cast.align_mul = nir_intrinsic_align_mul(intr);