AlexIndustrial/mesa
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

nir: Drop "SSA" from NIR language

Browse Source 
Everything is SSA now.

   sed -e 's/nir_ssa_def/nir_def/g' \
       -e 's/nir_ssa_undef/nir_undef/g' \
       -e 's/nir_ssa_scalar/nir_scalar/g' \
       -e 's/nir_src_rewrite_ssa/nir_src_rewrite/g' \
       -e 's/nir_gather_ssa_types/nir_gather_types/g' \
       -i $(git grep -l nir | grep -v relnotes)

   git mv src/compiler/nir/nir_gather_ssa_types.c \
          src/compiler/nir/nir_gather_types.c

   ninja -C build/ clang-format
   cd src/compiler/nir && find *.c *.h -type f -exec clang-format -i \{} \;

Signed-off-by: Alyssa Rosenzweig <alyssa@rosenzweig.io>
Acked-by: Faith Ekstrand <faith.ekstrand@collabora.com>
Acked-by: Emma Anholt <emma@anholt.net>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/24585>
This commit is contained in:
Alyssa Rosenzweig
2023-08-12 16:17:15 -04:00
parent 777d336b1f
commit 09d31922de
492 changed files with 10408 additions and 10455 deletions
Download Patch File Download Diff File Expand all files Collapse all files
src/microsoft/compiler/dxil_nir.c
+110 -110
View File
@@ -39,9 +39,9 @@ cl_type_size_align(const struct glsl_type *type, unsigned *size,
*align = glsl_get_cl_alignment(type);
}
static nir_ssa_def *
static nir_def *
load_comps_to_vec(nir_builder *b, unsigned src_bit_size,
nir_ssa_def **src_comps, unsigned num_src_comps,
nir_def **src_comps, unsigned num_src_comps,
unsigned dst_bit_size)
{
if (src_bit_size == dst_bit_size)
@@ -51,14 +51,14 @@ load_comps_to_vec(nir_builder *b, unsigned src_bit_size,
unsigned num_dst_comps = DIV_ROUND_UP(num_src_comps * src_bit_size, dst_bit_size);
unsigned comps_per_dst = dst_bit_size / src_bit_size;
nir_ssa_def *dst_comps[4];
nir_def *dst_comps[4];
for (unsigned i = 0; i < num_dst_comps; i++) {
unsigned src_offs = i * comps_per_dst;
dst_comps[i] = nir_u2uN(b, src_comps[src_offs], dst_bit_size);
for (unsigned j = 1; j < comps_per_dst && src_offs + j < num_src_comps; j++) {
nir_ssa_def *tmp = nir_ishl_imm(b, nir_u2uN(b, src_comps[src_offs + j], dst_bit_size),
nir_def *tmp = nir_ishl_imm(b, nir_u2uN(b, src_comps[src_offs + j], dst_bit_size),
j * src_bit_size);
dst_comps[i] = nir_ior(b, dst_comps[i], tmp);
}
@@ -76,14 +76,14 @@ lower_32b_offset_load(nir_builder *b, nir_intrinsic_instr *intr, nir_variable *v
b->cursor = nir_before_instr(&intr->instr);
nir_ssa_def *offset = intr->src[0].ssa;
nir_def *offset = intr->src[0].ssa;
if (intr->intrinsic == nir_intrinsic_load_shared)
offset = nir_iadd_imm(b, offset, nir_intrinsic_base(intr));
else
offset = nir_u2u32(b, offset);
nir_ssa_def *index = nir_ushr_imm(b, offset, 2);
nir_ssa_def *comps[NIR_MAX_VEC_COMPONENTS];
nir_ssa_def *comps_32bit[NIR_MAX_VEC_COMPONENTS * 2];
nir_def *index = nir_ushr_imm(b, offset, 2);
nir_def *comps[NIR_MAX_VEC_COMPONENTS];
nir_def *comps_32bit[NIR_MAX_VEC_COMPONENTS * 2];
/* We need to split loads in 32-bit accesses because the buffer
* is an i32 array and DXIL does not support type casts.
@@ -96,40 +96,40 @@ lower_32b_offset_load(nir_builder *b, nir_intrinsic_instr *intr, nir_variable *v
for (unsigned i = 0; i < num_32bit_comps; i += num_comps_per_pass) {
unsigned num_vec32_comps = MIN2(num_32bit_comps - i, 4);
unsigned num_dest_comps = num_vec32_comps * 32 / bit_size;
nir_ssa_def *vec32 = nir_vec(b, &comps_32bit[i], num_vec32_comps);
nir_def *vec32 = nir_vec(b, &comps_32bit[i], num_vec32_comps);
/* If we have 16 bits or less to load we need to adjust the u32 value so
* we can always extract the LSB.
*/
if (num_bits <= 16) {
nir_ssa_def *shift =
nir_def *shift =
nir_imul_imm(b, nir_iand_imm(b, offset, 3), 8);
vec32 = nir_ushr(b, vec32, shift);
}
/* And now comes the pack/unpack step to match the original type. */
unsigned dest_index = i * 32 / bit_size;
nir_ssa_def *temp_vec = nir_extract_bits(b, &vec32, 1, 0, num_dest_comps, bit_size);
nir_def *temp_vec = nir_extract_bits(b, &vec32, 1, 0, num_dest_comps, bit_size);
for (unsigned comp = 0; comp < num_dest_comps; ++comp, ++dest_index)
comps[dest_index] = nir_channel(b, temp_vec, comp);
}
nir_ssa_def *result = nir_vec(b, comps, num_components);
nir_ssa_def_rewrite_uses(&intr->dest.ssa, result);
nir_def *result = nir_vec(b, comps, num_components);
nir_def_rewrite_uses(&intr->dest.ssa, result);
nir_instr_remove(&intr->instr);
return true;
}
static void
lower_masked_store_vec32(nir_builder *b, nir_ssa_def *offset, nir_ssa_def *index,
nir_ssa_def *vec32, unsigned num_bits, nir_variable *var, unsigned alignment)
lower_masked_store_vec32(nir_builder *b, nir_def *offset, nir_def *index,
nir_def *vec32, unsigned num_bits, nir_variable *var, unsigned alignment)
{
nir_ssa_def *mask = nir_imm_int(b, (1 << num_bits) - 1);
nir_def *mask = nir_imm_int(b, (1 << num_bits) - 1);
/* If we have small alignments, we need to place them correctly in the u32 component. */
if (alignment <= 2) {
nir_ssa_def *shift =
nir_def *shift =
nir_imul_imm(b, nir_iand_imm(b, offset, 3), 8);
vec32 = nir_ishl(b, vec32, shift);
@@ -143,9 +143,9 @@ lower_masked_store_vec32(nir_builder *b, nir_ssa_def *offset, nir_ssa_def *index
nir_deref_atomic(b, 32, &deref->dest.ssa, vec32, .atomic_op = nir_atomic_op_ior);
} else {
/* For scratch, since we don't need atomics, just generate the read-modify-write in NIR */
nir_ssa_def *load = nir_load_array_var(b, var, index);
nir_def *load = nir_load_array_var(b, var, index);
nir_ssa_def *new_val = nir_ior(b, vec32,
nir_def *new_val = nir_ior(b, vec32,
nir_iand(b,
nir_inot(b, mask),
load));
@@ -163,12 +163,12 @@ lower_32b_offset_store(nir_builder *b, nir_intrinsic_instr *intr, nir_variable *
b->cursor = nir_before_instr(&intr->instr);
nir_ssa_def *offset = intr->src[1].ssa;
nir_def *offset = intr->src[1].ssa;
if (intr->intrinsic == nir_intrinsic_store_shared)
offset = nir_iadd_imm(b, offset, nir_intrinsic_base(intr));
else
offset = nir_u2u32(b, offset);
nir_ssa_def *comps[NIR_MAX_VEC_COMPONENTS];
nir_def *comps[NIR_MAX_VEC_COMPONENTS];
unsigned comp_idx = 0;
for (unsigned i = 0; i < num_components; i++)
@@ -178,10 +178,10 @@ lower_32b_offset_store(nir_builder *b, nir_intrinsic_instr *intr, nir_variable *
for (unsigned i = 0; i < num_bits; i += step) {
/* For each 4byte chunk (or smaller) we generate a 32bit scalar store. */
unsigned substore_num_bits = MIN2(num_bits - i, step);
nir_ssa_def *local_offset = nir_iadd_imm(b, offset, i / 8);
nir_ssa_def *vec32 = load_comps_to_vec(b, bit_size, &comps[comp_idx],
nir_def *local_offset = nir_iadd_imm(b, offset, i / 8);
nir_def *vec32 = load_comps_to_vec(b, bit_size, &comps[comp_idx],
substore_num_bits / bit_size, 32);
nir_ssa_def *index = nir_ushr_imm(b, local_offset, 2);
nir_def *index = nir_ushr_imm(b, local_offset, 2);
/* For anything less than 32bits we need to use the masked version of the
* intrinsic to preserve data living in the same 32bit slot. */
@@ -272,7 +272,7 @@ dxil_nir_lower_constant_to_temp(nir_shader *nir)
deref->dest.ssa.bit_size = 32;
if (deref->deref_type == nir_deref_type_array) {
b.cursor = nir_before_instr(instr);
nir_src_rewrite_ssa(&deref->arr.index, nir_u2u32(&b, deref->arr.index.ssa));
nir_src_rewrite(&deref->arr.index, nir_u2u32(&b, deref->arr.index.ssa));
}
}
}
@@ -319,12 +319,12 @@ flatten_var_arrays(nir_builder *b, nir_instr *instr, void *data)
assert(path.path[0]->deref_type == nir_deref_type_var);
b->cursor = nir_before_instr(&path.path[0]->instr);
nir_deref_instr *new_var_deref = nir_build_deref_var(b, var);
nir_ssa_def *index = NULL;
nir_def *index = NULL;
for (unsigned level = 1; path.path[level]; ++level) {
nir_deref_instr *arr_deref = path.path[level];
assert(arr_deref->deref_type == nir_deref_type_array);
b->cursor = nir_before_instr(&arr_deref->instr);
nir_ssa_def *val = nir_imul_imm(b, arr_deref->arr.index.ssa,
nir_def *val = nir_imul_imm(b, arr_deref->arr.index.ssa,
glsl_get_component_slots(arr_deref->type));
if (index) {
index = nir_iadd(b, index, val);
@@ -337,25 +337,25 @@ flatten_var_arrays(nir_builder *b, nir_instr *instr, void *data)
if (vector_comps > 1) {
b->cursor = nir_before_instr(instr);
if (intr->intrinsic == nir_intrinsic_load_deref) {
nir_ssa_def *components[NIR_MAX_VEC_COMPONENTS];
nir_def *components[NIR_MAX_VEC_COMPONENTS];
for (unsigned i = 0; i < vector_comps; ++i) {
nir_ssa_def *final_index = index ? nir_iadd_imm(b, index, i) : nir_imm_int(b, i);
nir_def *final_index = index ? nir_iadd_imm(b, index, i) : nir_imm_int(b, i);
nir_deref_instr *comp_deref = nir_build_deref_array(b, new_var_deref, final_index);
components[i] = nir_load_deref(b, comp_deref);
}
nir_ssa_def_rewrite_uses(&intr->dest.ssa, nir_vec(b, components, vector_comps));
nir_def_rewrite_uses(&intr->dest.ssa, nir_vec(b, components, vector_comps));
} else if (intr->intrinsic == nir_intrinsic_store_deref) {
for (unsigned i = 0; i < vector_comps; ++i) {
if (((1 << i) & nir_intrinsic_write_mask(intr)) == 0)
continue;
nir_ssa_def *final_index = index ? nir_iadd_imm(b, index, i) : nir_imm_int(b, i);
nir_def *final_index = index ? nir_iadd_imm(b, index, i) : nir_imm_int(b, i);
nir_deref_instr *comp_deref = nir_build_deref_array(b, new_var_deref, final_index);
nir_store_deref(b, comp_deref, nir_channel(b, intr->src[1].ssa, i), 1);
}
}
nir_instr_remove(instr);
} else {
nir_src_rewrite_ssa(&intr->src[0], &nir_build_deref_array(b, new_var_deref, index)->dest.ssa);
nir_src_rewrite(&intr->src[0], &nir_build_deref_array(b, new_var_deref, index)->dest.ssa);
}
nir_deref_path_finish(&path);
@@ -462,13 +462,13 @@ lower_deref_bit_size(nir_builder *b, nir_instr *instr, void *data)
if (intr->intrinsic == nir_intrinsic_load_deref) {
intr->dest.ssa.bit_size = glsl_get_bit_size(var_scalar_type);
b->cursor = nir_after_instr(instr);
nir_ssa_def *downcast = nir_type_convert(b, &intr->dest.ssa, new_type, old_type, nir_rounding_mode_undef);
nir_ssa_def_rewrite_uses_after(&intr->dest.ssa, downcast, downcast->parent_instr);
nir_def *downcast = nir_type_convert(b, &intr->dest.ssa, new_type, old_type, nir_rounding_mode_undef);
nir_def_rewrite_uses_after(&intr->dest.ssa, downcast, downcast->parent_instr);
}
else {
b->cursor = nir_before_instr(instr);
nir_ssa_def *upcast = nir_type_convert(b, intr->src[1].ssa, old_type, new_type, nir_rounding_mode_undef);
nir_src_rewrite_ssa(&intr->src[1], upcast);
nir_def *upcast = nir_type_convert(b, intr->src[1].ssa, old_type, new_type, nir_rounding_mode_undef);
nir_src_rewrite(&intr->src[1], upcast);
}
while (deref->deref_type == nir_deref_type_array) {
@@ -488,12 +488,12 @@ lower_deref_bit_size(nir_builder *b, nir_instr *instr, void *data)
nir_iadd_imm(b, deref->arr.index.ssa, 1));
b->cursor = nir_before_instr(instr);
if (intr->intrinsic == nir_intrinsic_load_deref) {
nir_ssa_def *src1 = nir_load_deref(b, deref);
nir_ssa_def *src2 = nir_load_deref(b, deref2);
nir_ssa_def_rewrite_uses(&intr->dest.ssa, nir_pack_64_2x32_split(b, src1, src2));
nir_def *src1 = nir_load_deref(b, deref);
nir_def *src2 = nir_load_deref(b, deref2);
nir_def_rewrite_uses(&intr->dest.ssa, nir_pack_64_2x32_split(b, src1, src2));
} else {
nir_ssa_def *src1 = nir_unpack_64_2x32_split_x(b, intr->src[1].ssa);
nir_ssa_def *src2 = nir_unpack_64_2x32_split_y(b, intr->src[1].ssa);
nir_def *src1 = nir_unpack_64_2x32_split_x(b, intr->src[1].ssa);
nir_def *src2 = nir_unpack_64_2x32_split_y(b, intr->src[1].ssa);
nir_store_deref(b, deref, src1, 1);
nir_store_deref(b, deref, src2, 1);
}
@@ -619,12 +619,12 @@ lower_shared_atomic(nir_builder *b, nir_intrinsic_instr *intr, nir_variable *var
{
b->cursor = nir_before_instr(&intr->instr);
nir_ssa_def *offset =
nir_def *offset =
nir_iadd_imm(b, intr->src[0].ssa, nir_intrinsic_base(intr));
nir_ssa_def *index = nir_ushr_imm(b, offset, 2);
nir_def *index = nir_ushr_imm(b, offset, 2);
nir_deref_instr *deref = nir_build_deref_array(b, nir_build_deref_var(b, var), index);
nir_ssa_def *result;
nir_def *result;
if (intr->intrinsic == nir_intrinsic_shared_atomic_swap)
result = nir_deref_atomic_swap(b, 32, &deref->dest.ssa, intr->src[1].ssa, intr->src[2].ssa,
.atomic_op = nir_intrinsic_atomic_op(intr));
@@ -632,7 +632,7 @@ lower_shared_atomic(nir_builder *b, nir_intrinsic_instr *intr, nir_variable *var
result = nir_deref_atomic(b, 32, &deref->dest.ssa, intr->src[1].ssa,
.atomic_op = nir_intrinsic_atomic_op(intr));
nir_ssa_def_rewrite_uses(&intr->dest.ssa, result);
nir_def_rewrite_uses(&intr->dest.ssa, result);
nir_instr_remove(&intr->instr);
return true;
}
@@ -713,11 +713,11 @@ lower_deref_ssbo(nir_builder *b, nir_deref_instr *deref)
/* We turn all deref_var into deref_cast and build a pointer value based on
* the var binding which encodes the UAV id.
*/
nir_ssa_def *ptr = nir_imm_int64(b, (uint64_t)var->data.binding << 32);
nir_def *ptr = nir_imm_int64(b, (uint64_t)var->data.binding << 32);
nir_deref_instr *deref_cast =
nir_build_deref_cast(b, ptr, nir_var_mem_ssbo, deref->type,
glsl_get_explicit_stride(var->type));
nir_ssa_def_rewrite_uses(&deref->dest.ssa,
nir_def_rewrite_uses(&deref->dest.ssa,
&deref_cast->dest.ssa);
nir_instr_remove(&deref->instr);
@@ -781,7 +781,7 @@ lower_alu_deref_srcs(nir_builder *b, nir_alu_instr *alu)
if (root_deref->deref_type != nir_deref_type_cast)
continue;
nir_ssa_def *ptr =
nir_def *ptr =
nir_iadd(b, root_deref->parent.ssa,
nir_build_deref_offset(b, deref, cl_type_size_align));
nir_instr_rewrite_src(&alu->instr, &alu->src[i].src, nir_src_for_ssa(ptr));
@@ -830,7 +830,7 @@ cast_phi(nir_builder *b, nir_phi_instr *phi, unsigned new_bit_size)
b->cursor = nir_after_instr_and_phis(src->src.ssa->parent_instr);
nir_ssa_def *cast = nir_u2uN(b, src->src.ssa, new_bit_size);
nir_def *cast = nir_u2uN(b, src->src.ssa, new_bit_size);
nir_phi_instr_add_src(lowered, src->pred, nir_src_for_ssa(cast));
}
@@ -842,9 +842,9 @@ cast_phi(nir_builder *b, nir_phi_instr *phi, unsigned new_bit_size)
nir_builder_instr_insert(b, &lowered->instr);
b->cursor = nir_after_phis(nir_cursor_current_block(b->cursor));
nir_ssa_def *result = nir_u2uN(b, &lowered->dest.ssa, old_bit_size);
nir_def *result = nir_u2uN(b, &lowered->dest.ssa, old_bit_size);
nir_ssa_def_rewrite_uses(&phi->dest.ssa, result);
nir_def_rewrite_uses(&phi->dest.ssa, result);
nir_instr_remove(&phi->instr);
}
@@ -1002,7 +1002,7 @@ dxil_nir_split_clip_cull_distance_instr(nir_builder *b,
new_intermediate_deref = nir_build_deref_array(b, new_intermediate_deref, parent->arr.index.ssa);
}
nir_deref_instr *new_array_deref = nir_build_deref_array(b, new_intermediate_deref, nir_imm_int(b, total_index % 4));
nir_ssa_def_rewrite_uses(&deref->dest.ssa, &new_array_deref->dest.ssa);
nir_def_rewrite_uses(&deref->dest.ssa, &new_array_deref->dest.ssa);
return true;
}
@@ -1048,10 +1048,10 @@ dxil_nir_lower_double_math_instr(nir_builder *b,
unsigned num_components = nir_op_infos[alu->op].input_sizes[i];
if (!num_components)
num_components = alu->dest.dest.ssa.num_components;
nir_ssa_def *components[NIR_MAX_VEC_COMPONENTS];
nir_def *components[NIR_MAX_VEC_COMPONENTS];
for (unsigned c = 0; c < num_components; ++c) {
nir_ssa_def *packed_double = nir_channel(b, alu->src[i].src.ssa, alu->src[i].swizzle[c]);
nir_ssa_def *unpacked_double = nir_unpack_64_2x32(b, packed_double);
nir_def *packed_double = nir_channel(b, alu->src[i].src.ssa, alu->src[i].swizzle[c]);
nir_def *unpacked_double = nir_unpack_64_2x32(b, packed_double);
components[c] = nir_pack_double_2x32_dxil(b, unpacked_double);
alu->src[i].swizzle[c] = c;
}
@@ -1063,14 +1063,14 @@ dxil_nir_lower_double_math_instr(nir_builder *b,
if (nir_alu_type_get_base_type(nir_op_infos[alu->op].output_type) == nir_type_float &&
alu->dest.dest.ssa.bit_size == 64) {
b->cursor = nir_after_instr(&alu->instr);
nir_ssa_def *components[NIR_MAX_VEC_COMPONENTS];
nir_def *components[NIR_MAX_VEC_COMPONENTS];
for (unsigned c = 0; c < alu->dest.dest.ssa.num_components; ++c) {
nir_ssa_def *packed_double = nir_channel(b, &alu->dest.dest.ssa, c);
nir_ssa_def *unpacked_double = nir_unpack_double_2x32_dxil(b, packed_double);
nir_def *packed_double = nir_channel(b, &alu->dest.dest.ssa, c);
nir_def *unpacked_double = nir_unpack_double_2x32_dxil(b, packed_double);
components[c] = nir_pack_64_2x32(b, unpacked_double);
}
nir_ssa_def *repacked_dvec = nir_vec(b, components, alu->dest.dest.ssa.num_components);
nir_ssa_def_rewrite_uses_after(&alu->dest.dest.ssa, repacked_dvec, repacked_dvec->parent_instr);
nir_def *repacked_dvec = nir_vec(b, components, alu->dest.dest.ssa.num_components);
nir_def_rewrite_uses_after(&alu->dest.dest.ssa, repacked_dvec, repacked_dvec->parent_instr);
progress = true;
}
@@ -1128,7 +1128,7 @@ lower_system_value_to_zero_filter(const nir_instr* instr, const void* cb_state)
return false;
}
static nir_ssa_def*
static nir_def*
lower_system_value_to_zero_instr(nir_builder* b, nir_instr* instr, void* _state)
{
return nir_imm_int(b, 0);
@@ -1156,8 +1156,8 @@ lower_load_local_group_size(nir_builder *b, nir_intrinsic_instr *intr)
nir_const_value_for_int(b->shader->info.workgroup_size[1], 32),
nir_const_value_for_int(b->shader->info.workgroup_size[2], 32)
};
nir_ssa_def *size = nir_build_imm(b, 3, 32, v);
nir_ssa_def_rewrite_uses(&intr->dest.ssa, size);
nir_def *size = nir_build_imm(b, 3, 32, v);
nir_def_rewrite_uses(&intr->dest.ssa, size);
nir_instr_remove(&intr->instr);
}
@@ -1423,7 +1423,7 @@ lower_sysval_to_load_input_impl(nir_builder *b, nir_instr *instr, void *data)
? 32 : intr->dest.ssa.bit_size;
b->cursor = nir_before_instr(instr);
nir_ssa_def *result = nir_load_input(b, intr->dest.ssa.num_components, bit_size, nir_imm_int(b, 0),
nir_def *result = nir_load_input(b, intr->dest.ssa.num_components, bit_size, nir_imm_int(b, 0),
.base = var->data.driver_location, .dest_type = dest_type);
/* The nir_type_uint32 is really a nir_type_bool32, but that type is very
@@ -1433,7 +1433,7 @@ lower_sysval_to_load_input_impl(nir_builder *b, nir_instr *instr, void *data)
if (sysval == SYSTEM_VALUE_FRONT_FACE)
result = nir_ine_imm(b, result, 0);
nir_ssa_def_rewrite_uses(&intr->dest.ssa, result);
nir_def_rewrite_uses(&intr->dest.ssa, result);
return true;
}
@@ -1610,14 +1610,14 @@ lower_ubo_array_one_to_static(struct nir_builder *b, nir_instr *inst,
// Indexing out of bounds on array of UBOs is considered undefined
// behavior. Therefore, we just hardcode all the index to 0.
uint8_t bit_size = index->dest.ssa.bit_size;
nir_ssa_def *zero = nir_imm_intN_t(b, 0, bit_size);
nir_ssa_def *dest =
nir_def *zero = nir_imm_intN_t(b, 0, bit_size);
nir_def *dest =
nir_vulkan_resource_index(b, index->num_components, bit_size, zero,
.desc_set = nir_intrinsic_desc_set(index),
.binding = nir_intrinsic_binding(index),
.desc_type = nir_intrinsic_desc_type(index));
nir_ssa_def_rewrite_uses(&index->dest.ssa, dest);
nir_def_rewrite_uses(&index->dest.ssa, dest);
return true;
}
@@ -1641,7 +1641,7 @@ is_fquantize2f16(const nir_instr *instr, const void *data)
return alu->op == nir_op_fquantize2f16;
}
static nir_ssa_def *
static nir_def *
lower_fquantize2f16(struct nir_builder *b, nir_instr *instr, void *data)
{
/*
@@ -1672,19 +1672,19 @@ lower_fquantize2f16(struct nir_builder *b, nir_instr *instr, void *data)
* return round(val);
*/
nir_alu_instr *alu = nir_instr_as_alu(instr);
nir_ssa_def *src =
nir_def *src =
nir_ssa_for_src(b, alu->src[0].src, nir_src_num_components(alu->src[0].src));
nir_ssa_def *neg_inf_cond =
nir_def *neg_inf_cond =
nir_flt_imm(b, src, -65504.0f);
nir_ssa_def *pos_inf_cond =
nir_def *pos_inf_cond =
nir_fgt_imm(b, src, 65504.0f);
nir_ssa_def *zero_cond =
nir_def *zero_cond =
nir_flt_imm(b, nir_fabs(b, src), ldexpf(1.0, -14));
nir_ssa_def *zero = nir_iand_imm(b, src, 1 << 31);
nir_ssa_def *round = nir_iand_imm(b, src, ~BITFIELD_MASK(13));
nir_def *zero = nir_iand_imm(b, src, 1 << 31);
nir_def *round = nir_iand_imm(b, src, ~BITFIELD_MASK(13));
nir_ssa_def *res =
nir_def *res =
nir_bcsel(b, neg_inf_cond, nir_imm_float(b, -INFINITY), round);
res = nir_bcsel(b, pos_inf_cond, nir_imm_float(b, INFINITY), res);
res = nir_bcsel(b, zero_cond, zero, res);
@@ -1847,14 +1847,14 @@ update_writes(struct nir_builder *b, nir_instr *instr, void *_state)
if (io.location != VARYING_SLOT_POS)
return false;
nir_ssa_def *src = intr->src[0].ssa;
nir_def *src = intr->src[0].ssa;
unsigned write_mask = nir_intrinsic_write_mask(intr);
if (src->num_components == 4 && write_mask == 0xf)
return false;
b->cursor = nir_before_instr(instr);
unsigned first_comp = nir_intrinsic_component(intr);
nir_ssa_def *channels[4] = { NULL, NULL, NULL, NULL };
nir_def *channels[4] = { NULL, NULL, NULL, NULL };
assert(first_comp + src->num_components <= ARRAY_SIZE(channels));
for (unsigned i = 0; i < src->num_components; ++i)
if (write_mask & (1 << i))
@@ -1894,7 +1894,7 @@ is_sample_pos(const nir_instr *instr, const void *_data)
return intr->intrinsic == nir_intrinsic_load_sample_pos;
}
static nir_ssa_def *
static nir_def *
lower_sample_pos(nir_builder *b, nir_instr *instr, void *_data)
{
return nir_load_sample_pos_from_id(b, 32, nir_load_sample_id(b));
@@ -1920,19 +1920,19 @@ lower_subgroup_id(nir_builder *b, nir_instr *instr, void *data)
b->shader->info.workgroup_size[2] == 1) {
/* When using Nx1x1 groups, use a simple stable algorithm
* which is almost guaranteed to be correct. */
nir_ssa_def *subgroup_id = nir_udiv(b, nir_load_local_invocation_index(b), nir_load_subgroup_size(b));
nir_ssa_def_rewrite_uses(&intr->dest.ssa, subgroup_id);
nir_def *subgroup_id = nir_udiv(b, nir_load_local_invocation_index(b), nir_load_subgroup_size(b));
nir_def_rewrite_uses(&intr->dest.ssa, subgroup_id);
return true;
}
nir_ssa_def **subgroup_id = (nir_ssa_def **)data;
nir_def **subgroup_id = (nir_def **)data;
if (*subgroup_id == NULL) {
nir_variable *subgroup_id_counter = nir_variable_create(b->shader, nir_var_mem_shared, glsl_uint_type(), "dxil_SubgroupID_counter");
nir_variable *subgroup_id_local = nir_local_variable_create(b->impl, glsl_uint_type(), "dxil_SubgroupID_local");
nir_store_var(b, subgroup_id_local, nir_imm_int(b, 0), 1);
nir_deref_instr *counter_deref = nir_build_deref_var(b, subgroup_id_counter);
nir_ssa_def *tid = nir_load_local_invocation_index(b);
nir_def *tid = nir_load_local_invocation_index(b);
nir_if *nif = nir_push_if(b, nir_ieq_imm(b, tid, 0));
nir_store_deref(b, counter_deref, nir_imm_int(b, 0), 1);
nir_pop_if(b, nif);
@@ -1944,22 +1944,22 @@ lower_subgroup_id(nir_builder *b, nir_instr *instr, void *data)
.memory_modes = nir_var_mem_shared);
nif = nir_push_if(b, nir_elect(b, 1));
nir_ssa_def *subgroup_id_first_thread = nir_deref_atomic(b, 32, &counter_deref->dest.ssa, nir_imm_int(b, 1),
nir_def *subgroup_id_first_thread = nir_deref_atomic(b, 32, &counter_deref->dest.ssa, nir_imm_int(b, 1),
.atomic_op = nir_atomic_op_iadd);
nir_store_var(b, subgroup_id_local, subgroup_id_first_thread, 1);
nir_pop_if(b, nif);
nir_ssa_def *subgroup_id_loaded = nir_load_var(b, subgroup_id_local);
nir_def *subgroup_id_loaded = nir_load_var(b, subgroup_id_local);
*subgroup_id = nir_read_first_invocation(b, subgroup_id_loaded);
}
nir_ssa_def_rewrite_uses(&intr->dest.ssa, *subgroup_id);
nir_def_rewrite_uses(&intr->dest.ssa, *subgroup_id);
return true;
}
bool
dxil_nir_lower_subgroup_id(nir_shader *s)
{
nir_ssa_def *subgroup_id = NULL;
nir_def *subgroup_id = NULL;
return nir_shader_instructions_pass(s, lower_subgroup_id, nir_metadata_none, &subgroup_id);
}
@@ -1973,14 +1973,14 @@ lower_num_subgroups(nir_builder *b, nir_instr *instr, void *data)
return false;
b->cursor = nir_before_instr(instr);
nir_ssa_def *subgroup_size = nir_load_subgroup_size(b);
nir_ssa_def *size_minus_one = nir_iadd_imm(b, subgroup_size, -1);
nir_ssa_def *workgroup_size_vec = nir_load_workgroup_size(b);
nir_ssa_def *workgroup_size = nir_imul(b, nir_channel(b, workgroup_size_vec, 0),
nir_def *subgroup_size = nir_load_subgroup_size(b);
nir_def *size_minus_one = nir_iadd_imm(b, subgroup_size, -1);
nir_def *workgroup_size_vec = nir_load_workgroup_size(b);
nir_def *workgroup_size = nir_imul(b, nir_channel(b, workgroup_size_vec, 0),
nir_imul(b, nir_channel(b, workgroup_size_vec, 1),
nir_channel(b, workgroup_size_vec, 2)));
nir_ssa_def *ret = nir_idiv(b, nir_iadd(b, workgroup_size, size_minus_one), subgroup_size);
nir_ssa_def_rewrite_uses(&intr->dest.ssa, ret);
nir_def *ret = nir_idiv(b, nir_iadd(b, workgroup_size, size_minus_one), subgroup_size);
nir_def_rewrite_uses(&intr->dest.ssa, ret);
return true;
}
@@ -2014,7 +2014,7 @@ static void
split_unaligned_load(nir_builder *b, nir_intrinsic_instr *intrin, unsigned alignment)
{
enum gl_access_qualifier access = nir_intrinsic_access(intrin);
nir_ssa_def *srcs[NIR_MAX_VEC_COMPONENTS * NIR_MAX_VEC_COMPONENTS * sizeof(int64_t) / 8];
nir_def *srcs[NIR_MAX_VEC_COMPONENTS * NIR_MAX_VEC_COMPONENTS * sizeof(int64_t) / 8];
unsigned comp_size = intrin->dest.ssa.bit_size / 8;
unsigned num_comps = intrin->dest.ssa.num_components;
@@ -2031,8 +2031,8 @@ split_unaligned_load(nir_builder *b, nir_intrinsic_instr *intrin, unsigned align
srcs[i] = nir_load_deref_with_access(b, elem, access);
}
nir_ssa_def *new_dest = nir_extract_bits(b, srcs, num_loads, 0, num_comps, intrin->dest.ssa.bit_size);
nir_ssa_def_rewrite_uses(&intrin->dest.ssa, new_dest);
nir_def *new_dest = nir_extract_bits(b, srcs, num_loads, 0, num_comps, intrin->dest.ssa.bit_size);
nir_def_rewrite_uses(&intrin->dest.ssa, new_dest);
nir_instr_remove(&intrin->instr);
}
@@ -2041,7 +2041,7 @@ split_unaligned_store(nir_builder *b, nir_intrinsic_instr *intrin, unsigned alig
{
enum gl_access_qualifier access = nir_intrinsic_access(intrin);
nir_ssa_def *value = intrin->src[1].ssa;
nir_def *value = intrin->src[1].ssa;
unsigned comp_size = value->bit_size / 8;
unsigned num_comps = value->num_components;
@@ -2054,7 +2054,7 @@ split_unaligned_store(nir_builder *b, nir_intrinsic_instr *intrin, unsigned alig
unsigned num_stores = DIV_ROUND_UP(comp_size * num_comps, alignment);
for (unsigned i = 0; i < num_stores; ++i) {
nir_ssa_def *substore_val = nir_extract_bits(b, &value, 1, i * alignment * 8, 1, alignment * 8);
nir_def *substore_val = nir_extract_bits(b, &value, 1, i * alignment * 8, 1, alignment * 8);
nir_deref_instr *elem = nir_build_deref_ptr_as_array(b, cast, nir_imm_intN_t(b, i, cast->dest.ssa.bit_size));
nir_store_deref_with_access(b, elem, substore_val, ~0, access);
}
@@ -2094,7 +2094,7 @@ dxil_nir_split_unaligned_loads_stores(nir_shader *shader, nir_variable_mode mode
if (alignment >= req_align)
continue;
nir_ssa_def *val;
nir_def *val;
if (intrin->intrinsic == nir_intrinsic_load_deref) {
val = &intrin->dest.ssa;
} else {
@@ -2133,9 +2133,9 @@ lower_inclusive_to_exclusive(nir_builder *b, nir_intrinsic_instr *intr)
intr->intrinsic = nir_intrinsic_exclusive_scan;
nir_intrinsic_set_reduction_op(intr, op);
nir_ssa_def *final_val = nir_build_alu2(b, nir_intrinsic_reduction_op(intr),
nir_def *final_val = nir_build_alu2(b, nir_intrinsic_reduction_op(intr),
&intr->dest.ssa, intr->src[0].ssa);
nir_ssa_def_rewrite_uses_after(&intr->dest.ssa, final_val, final_val->parent_instr);
nir_def_rewrite_uses_after(&intr->dest.ssa, final_val, final_val->parent_instr);
}
static bool
@@ -2166,9 +2166,9 @@ lower_subgroup_scan(nir_builder *b, nir_instr *instr, void *data)
b->cursor = nir_before_instr(instr);
nir_op op = nir_intrinsic_reduction_op(intr);
nir_ssa_def *subgroup_id = nir_load_subgroup_invocation(b);
nir_ssa_def *active_threads = nir_ballot(b, 4, 32, nir_imm_true(b));
nir_ssa_def *base_value;
nir_def *subgroup_id = nir_load_subgroup_invocation(b);
nir_def *active_threads = nir_ballot(b, 4, 32, nir_imm_true(b));
nir_def *base_value;
uint32_t bit_size = intr->dest.ssa.bit_size;
if (op == nir_op_iand || op == nir_op_umin)
base_value = nir_imm_intN_t(b, ~0ull, bit_size);
@@ -2191,12 +2191,12 @@ lower_subgroup_scan(nir_builder *b, nir_instr *instr, void *data)
nir_store_var(b, loop_counter_var, nir_imm_int(b, 0), 1);
nir_store_var(b, result_var, base_value, 1);
nir_loop *loop = nir_push_loop(b);
nir_ssa_def *loop_counter = nir_load_var(b, loop_counter_var);
nir_def *loop_counter = nir_load_var(b, loop_counter_var);
nir_if *nif = nir_push_if(b, intr->intrinsic == nir_intrinsic_inclusive_scan ?
nir_ige(b, subgroup_id, loop_counter) :
nir_ilt(b, loop_counter, subgroup_id));
nir_if *if_active_thread = nir_push_if(b, nir_ballot_bitfield_extract(b, 32, active_threads, loop_counter));
nir_ssa_def *result = nir_build_alu2(b, op,
nir_def *result = nir_build_alu2(b, op,
nir_load_var(b, result_var),
nir_read_invocation(b, intr->src[0].ssa, loop_counter));
nir_store_var(b, result_var, result, 1);
@@ -2208,7 +2208,7 @@ lower_subgroup_scan(nir_builder *b, nir_instr *instr, void *data)
nir_pop_loop(b, loop);
result = nir_load_var(b, result_var);
nir_ssa_def_rewrite_uses(&intr->dest.ssa, result);
nir_def_rewrite_uses(&intr->dest.ssa, result);
return true;
}
@@ -2237,9 +2237,9 @@ lower_load_face(nir_builder *b, nir_instr *instr, void *data)
b->cursor = nir_before_instr(&intr->instr);
nir_variable *var = data;
nir_ssa_def *load = nir_ine_imm(b, nir_load_var(b, var), 0);
nir_def *load = nir_ine_imm(b, nir_load_var(b, var), 0);
nir_ssa_def_rewrite_uses(&intr->dest.ssa, load);
nir_def_rewrite_uses(&intr->dest.ssa, load);
nir_instr_remove(instr);
return true;
}
@@ -2276,7 +2276,7 @@ move_consts(nir_builder *b, nir_instr *instr, void *data)
load_const->def.bit_size);
memcpy(new_load->value, load_const->value, sizeof(load_const->value[0]) * load_const->def.num_components);
nir_builder_instr_insert(b, &new_load->instr);
nir_src_rewrite_ssa(src, &new_load->def);
nir_src_rewrite(src, &new_load->def);
progress = true;
}
}
src/microsoft/compiler/dxil_nir_lower_int_cubemaps.c
+86 -86
View File
@@ -95,76 +95,76 @@ lower_int_cubmap_to_array_filter(const nir_instr *instr,
typedef struct {
bool image;
nir_ssa_def *rx;
nir_ssa_def *ry;
nir_ssa_def *rz;
nir_ssa_def *arx;
nir_ssa_def *ary;
nir_ssa_def *arz;
nir_ssa_def *array;
nir_def *rx;
nir_def *ry;
nir_def *rz;
nir_def *arx;
nir_def *ary;
nir_def *arz;
nir_def *array;
} coord_t;
/* This is taken from from sp_tex_sample:convert_cube */
static nir_ssa_def *
static nir_def *
evaluate_face_x(nir_builder *b, coord_t *coord)
{
nir_ssa_def *sign = nir_fsign(b, coord->rx);
nir_ssa_def *positive = nir_fge_imm(b, coord->rx, 0.0);
nir_ssa_def *ima = nir_fdiv(b, nir_imm_float(b, -0.5), coord->arx);
nir_def *sign = nir_fsign(b, coord->rx);
nir_def *positive = nir_fge_imm(b, coord->rx, 0.0);
nir_def *ima = nir_fdiv(b, nir_imm_float(b, -0.5), coord->arx);
nir_ssa_def *x = nir_fadd_imm(b, nir_fmul(b, nir_fmul(b, sign, ima), coord->rz), 0.5);
nir_ssa_def *y = nir_fadd_imm(b, nir_fmul(b, ima, coord->ry), 0.5);
nir_ssa_def *face = nir_bcsel(b, positive, nir_imm_float(b, 0.0), nir_imm_float(b, 1.0));
nir_def *x = nir_fadd_imm(b, nir_fmul(b, nir_fmul(b, sign, ima), coord->rz), 0.5);
nir_def *y = nir_fadd_imm(b, nir_fmul(b, ima, coord->ry), 0.5);
nir_def *face = nir_bcsel(b, positive, nir_imm_float(b, 0.0), nir_imm_float(b, 1.0));
if (coord->array)
face = nir_fadd(b, face, coord->array);
return coord->image ?
nir_vec4(b, x,y, face, nir_ssa_undef(b, 1, 32)) :
nir_vec4(b, x,y, face, nir_undef(b, 1, 32)) :
nir_vec3(b, x,y, face);
}
static nir_ssa_def *
static nir_def *
evaluate_face_y(nir_builder *b, coord_t *coord)
{
nir_ssa_def *sign = nir_fsign(b, coord->ry);
nir_ssa_def *positive = nir_fge_imm(b, coord->ry, 0.0);
nir_ssa_def *ima = nir_fdiv(b, nir_imm_float(b, 0.5), coord->ary);
nir_def *sign = nir_fsign(b, coord->ry);
nir_def *positive = nir_fge_imm(b, coord->ry, 0.0);
nir_def *ima = nir_fdiv(b, nir_imm_float(b, 0.5), coord->ary);
nir_ssa_def *x = nir_fadd_imm(b, nir_fmul(b, ima, coord->rx), 0.5);
nir_ssa_def *y = nir_fadd_imm(b, nir_fmul(b, nir_fmul(b, sign, ima), coord->rz), 0.5);
nir_ssa_def *face = nir_bcsel(b, positive, nir_imm_float(b, 2.0), nir_imm_float(b, 3.0));
nir_def *x = nir_fadd_imm(b, nir_fmul(b, ima, coord->rx), 0.5);
nir_def *y = nir_fadd_imm(b, nir_fmul(b, nir_fmul(b, sign, ima), coord->rz), 0.5);
nir_def *face = nir_bcsel(b, positive, nir_imm_float(b, 2.0), nir_imm_float(b, 3.0));
if (coord->array)
face = nir_fadd(b, face, coord->array);
return coord->image ?
nir_vec4(b, x,y, face, nir_ssa_undef(b, 1, 32)) :
nir_vec4(b, x,y, face, nir_undef(b, 1, 32)) :
nir_vec3(b, x,y, face);
}
static nir_ssa_def *
static nir_def *
evaluate_face_z(nir_builder *b, coord_t *coord)
{
nir_ssa_def *sign = nir_fsign(b, coord->rz);
nir_ssa_def *positive = nir_fge_imm(b, coord->rz, 0.0);
nir_ssa_def *ima = nir_fdiv(b, nir_imm_float(b, -0.5), coord->arz);
nir_def *sign = nir_fsign(b, coord->rz);
nir_def *positive = nir_fge_imm(b, coord->rz, 0.0);
nir_def *ima = nir_fdiv(b, nir_imm_float(b, -0.5), coord->arz);
nir_ssa_def *x = nir_fadd_imm(b, nir_fmul(b, nir_fmul(b, sign, ima), nir_fneg(b, coord->rx)), 0.5);
nir_ssa_def *y = nir_fadd_imm(b, nir_fmul(b, ima, coord->ry), 0.5);
nir_ssa_def *face = nir_bcsel(b, positive, nir_imm_float(b, 4.0), nir_imm_float(b, 5.0));
nir_def *x = nir_fadd_imm(b, nir_fmul(b, nir_fmul(b, sign, ima), nir_fneg(b, coord->rx)), 0.5);
nir_def *y = nir_fadd_imm(b, nir_fmul(b, ima, coord->ry), 0.5);
nir_def *face = nir_bcsel(b, positive, nir_imm_float(b, 4.0), nir_imm_float(b, 5.0));
if (coord->array)
face = nir_fadd(b, face, coord->array);
return coord->image ?
nir_vec4(b, x,y, face, nir_ssa_undef(b, 1, 32)) :
nir_vec4(b, x,y, face, nir_undef(b, 1, 32)) :
nir_vec3(b, x,y, face);
}
static nir_ssa_def *
create_array_tex_from_cube_tex(nir_builder *b, nir_tex_instr *tex, nir_ssa_def *coord, nir_texop op)
static nir_def *
create_array_tex_from_cube_tex(nir_builder *b, nir_tex_instr *tex, nir_def *coord, nir_texop op)
{
nir_tex_instr *array_tex;
@@ -194,8 +194,8 @@ create_array_tex_from_cube_tex(nir_builder *b, nir_tex_instr *tex, nir_ssa_def *
return &array_tex->dest.ssa;
}
static nir_ssa_def *
handle_cube_edge(nir_builder *b, nir_ssa_def *x, nir_ssa_def *y, nir_ssa_def *face, nir_ssa_def *array_slice_cube_base, nir_ssa_def *tex_size)
static nir_def *
handle_cube_edge(nir_builder *b, nir_def *x, nir_def *y, nir_def *face, nir_def *array_slice_cube_base, nir_def *tex_size)
{
enum cube_remap
{
@@ -263,22 +263,22 @@ handle_cube_edge(nir_builder *b, nir_ssa_def *x, nir_ssa_def *y, nir_ssa_def *fa
cube_remap_pos_y
};
nir_ssa_def *zero = nir_imm_int(b, 0);
nir_def *zero = nir_imm_int(b, 0);
/* Doesn't matter since the texture is square */
tex_size = nir_channel(b, tex_size, 0);
nir_ssa_def *x_on = nir_iand(b, nir_ige(b, x, zero), nir_ige(b, tex_size, x));
nir_ssa_def *y_on = nir_iand(b, nir_ige(b, y, zero), nir_ige(b, tex_size, y));
nir_ssa_def *one_on = nir_ixor(b, x_on, y_on);
nir_def *x_on = nir_iand(b, nir_ige(b, x, zero), nir_ige(b, tex_size, x));
nir_def *y_on = nir_iand(b, nir_ige(b, y, zero), nir_ige(b, tex_size, y));
nir_def *one_on = nir_ixor(b, x_on, y_on);
/* If the sample did not fall off the face in either dimension, then set output = input */
nir_ssa_def *x_result = x;
nir_ssa_def *y_result = y;
nir_ssa_def *face_result = face;
nir_def *x_result = x;
nir_def *y_result = y;
nir_def *face_result = face;
/* otherwise, if the sample fell off the face in either the X or the Y direction, remap to the new face */
nir_ssa_def *remap_predicates[4] =
nir_def *remap_predicates[4] =
{
nir_iand(b, one_on, nir_ilt(b, x, zero)),
nir_iand(b, one_on, nir_ilt(b, tex_size, x)),
@@ -286,7 +286,7 @@ handle_cube_edge(nir_builder *b, nir_ssa_def *x, nir_ssa_def *y, nir_ssa_def *fa
nir_iand(b, one_on, nir_ilt(b, tex_size, y)),
};
nir_ssa_def *remap_array[cube_remap_size];
nir_def *remap_array[cube_remap_size];
remap_array[cube_remap_zero] = zero;
remap_array[cube_remap_x] = x;
@@ -301,7 +301,7 @@ handle_cube_edge(nir_builder *b, nir_ssa_def *x, nir_ssa_def *y, nir_ssa_def *fa
/* For each possible original face */
for (unsigned j = 0; j < 6; j++) {
nir_ssa_def *predicate = nir_iand(b, remap_predicates[i], nir_ieq_imm(b, face, j));
nir_def *predicate = nir_iand(b, remap_predicates[i], nir_ieq_imm(b, face, j));
x_result = nir_bcsel(b, predicate, remap_array[remap_table[j].remap_x], x_result);
y_result = nir_bcsel(b, predicate, remap_array[remap_table[j].remap_y], y_result);
@@ -312,48 +312,48 @@ handle_cube_edge(nir_builder *b, nir_ssa_def *x, nir_ssa_def *y, nir_ssa_def *fa
return nir_vec3(b, x_result, y_result, nir_iadd(b, face_result, array_slice_cube_base));
}
static nir_ssa_def *
handle_cube_gather(nir_builder *b, nir_tex_instr *tex, nir_ssa_def *coord)
static nir_def *
handle_cube_gather(nir_builder *b, nir_tex_instr *tex, nir_def *coord)
{
nir_ssa_def *tex_size = nir_get_texture_size(b, tex);
nir_def *tex_size = nir_get_texture_size(b, tex);
/* nir_get_texture_size puts the cursor before the tex op */
b->cursor = nir_after_instr(coord->parent_instr);
nir_ssa_def *const_05 = nir_imm_float(b, 0.5f);
nir_ssa_def *texel_coords = nir_fmul(b, nir_trim_vector(b, coord, 2),
nir_def *const_05 = nir_imm_float(b, 0.5f);
nir_def *texel_coords = nir_fmul(b, nir_trim_vector(b, coord, 2),
nir_i2f32(b, nir_trim_vector(b, tex_size, 2)));
nir_ssa_def *x_orig = nir_channel(b, texel_coords, 0);
nir_ssa_def *y_orig = nir_channel(b, texel_coords, 1);
nir_def *x_orig = nir_channel(b, texel_coords, 0);
nir_def *y_orig = nir_channel(b, texel_coords, 1);
nir_ssa_def *x_pos = nir_f2i32(b, nir_fadd(b, x_orig, const_05));
nir_ssa_def *x_neg = nir_f2i32(b, nir_fsub(b, x_orig, const_05));
nir_ssa_def *y_pos = nir_f2i32(b, nir_fadd(b, y_orig, const_05));
nir_ssa_def *y_neg = nir_f2i32(b, nir_fsub(b, y_orig, const_05));
nir_ssa_def *coords[4][2] = {
nir_def *x_pos = nir_f2i32(b, nir_fadd(b, x_orig, const_05));
nir_def *x_neg = nir_f2i32(b, nir_fsub(b, x_orig, const_05));
nir_def *y_pos = nir_f2i32(b, nir_fadd(b, y_orig, const_05));
nir_def *y_neg = nir_f2i32(b, nir_fsub(b, y_orig, const_05));
nir_def *coords[4][2] = {
{ x_neg, y_pos },
{ x_pos, y_pos },
{ x_pos, y_neg },
{ x_neg, y_neg },
};
nir_ssa_def *array_slice_2d = nir_f2i32(b, nir_channel(b, coord, 2));
nir_ssa_def *face = nir_imod_imm(b, array_slice_2d, 6);
nir_ssa_def *array_slice_cube_base = nir_isub(b, array_slice_2d, face);
nir_def *array_slice_2d = nir_f2i32(b, nir_channel(b, coord, 2));
nir_def *face = nir_imod_imm(b, array_slice_2d, 6);
nir_def *array_slice_cube_base = nir_isub(b, array_slice_2d, face);
nir_ssa_def *channels[4];
nir_def *channels[4];
for (unsigned i = 0; i < 4; ++i) {
nir_ssa_def *final_coord = handle_cube_edge(b, coords[i][0], coords[i][1], face, array_slice_cube_base, tex_size);
nir_ssa_def *sampled_val = create_array_tex_from_cube_tex(b, tex, final_coord, nir_texop_txf);
nir_def *final_coord = handle_cube_edge(b, coords[i][0], coords[i][1], face, array_slice_cube_base, tex_size);
nir_def *sampled_val = create_array_tex_from_cube_tex(b, tex, final_coord, nir_texop_txf);
channels[i] = nir_channel(b, sampled_val, tex->component);
}
return nir_vec(b, channels, 4);
}
static nir_ssa_def *
lower_cube_coords(nir_builder *b, nir_ssa_def *coord, bool is_array, bool is_image)
static nir_def *
lower_cube_coords(nir_builder *b, nir_def *coord, bool is_array, bool is_image)
{
coord_t coords;
coords.image = is_image;
@@ -367,43 +367,43 @@ lower_cube_coords(nir_builder *b, nir_ssa_def *coord, bool is_array, bool is_ima
if (is_array)
coords.array = nir_fmul_imm(b, nir_channel(b, coord, 3), 6.0f);
nir_ssa_def *use_face_x = nir_iand(b,
nir_def *use_face_x = nir_iand(b,
nir_fge(b, coords.arx, coords.ary),
nir_fge(b, coords.arx, coords.arz));
nir_if *use_face_x_if = nir_push_if(b, use_face_x);
nir_ssa_def *face_x_coord = evaluate_face_x(b, &coords);
nir_def *face_x_coord = evaluate_face_x(b, &coords);
nir_if *use_face_x_else = nir_push_else(b, use_face_x_if);
nir_ssa_def *use_face_y = nir_iand(b,
nir_def *use_face_y = nir_iand(b,
nir_fge(b, coords.ary, coords.arx),
nir_fge(b, coords.ary, coords.arz));
nir_if *use_face_y_if = nir_push_if(b, use_face_y);
nir_ssa_def *face_y_coord = evaluate_face_y(b, &coords);
nir_def *face_y_coord = evaluate_face_y(b, &coords);
nir_if *use_face_y_else = nir_push_else(b, use_face_y_if);
nir_ssa_def *face_z_coord = evaluate_face_z(b, &coords);
nir_def *face_z_coord = evaluate_face_z(b, &coords);
nir_pop_if(b, use_face_y_else);
nir_ssa_def *face_y_or_z_coord = nir_if_phi(b, face_y_coord, face_z_coord);
nir_def *face_y_or_z_coord = nir_if_phi(b, face_y_coord, face_z_coord);
nir_pop_if(b, use_face_x_else);
// This contains in xy the normalized sample coordinates, and in z the face index
nir_ssa_def *coord_and_face = nir_if_phi(b, face_x_coord, face_y_or_z_coord);
nir_def *coord_and_face = nir_if_phi(b, face_x_coord, face_y_or_z_coord);
return coord_and_face;
}
static nir_ssa_def *
static nir_def *
lower_cube_sample(nir_builder *b, nir_tex_instr *tex)
{
int coord_index = nir_tex_instr_src_index(tex, nir_tex_src_coord);
assert(coord_index >= 0);
/* Evaluate the face and the xy coordinates for a 2D tex op */
nir_ssa_def *coord = tex->src[coord_index].src.ssa;
nir_ssa_def *coord_and_face = lower_cube_coords(b, coord, tex->is_array, false);
nir_def *coord = tex->src[coord_index].src.ssa;
nir_def *coord_and_face = lower_cube_coords(b, coord, tex->is_array, false);
if (tex->op == nir_texop_tg4)
return handle_cube_gather(b, tex, coord_and_face);
@@ -411,7 +411,7 @@ lower_cube_sample(nir_builder *b, nir_tex_instr *tex)
return create_array_tex_from_cube_tex(b, tex, coord_and_face, tex->op);
}
static nir_ssa_def *
static nir_def *
lower_cube_image_load_store_atomic(nir_builder *b, nir_intrinsic_instr *intr)
{
b->cursor = nir_before_instr(&intr->instr);
@@ -421,29 +421,29 @@ lower_cube_image_load_store_atomic(nir_builder *b, nir_intrinsic_instr *intr)
return NIR_LOWER_INSTR_PROGRESS;
}
static nir_ssa_def *
static nir_def *
lower_cube_txs(nir_builder *b, nir_tex_instr *tex)
{
b->cursor = nir_after_instr(&tex->instr);
if (!tex->is_array)
return nir_trim_vector(b, &tex->dest.ssa, 2);
nir_ssa_def *array_dim = nir_channel(b, &tex->dest.ssa, 2);
nir_ssa_def *cube_array_dim = nir_idiv(b, array_dim, nir_imm_int(b, 6));
nir_def *array_dim = nir_channel(b, &tex->dest.ssa, 2);
nir_def *cube_array_dim = nir_idiv(b, array_dim, nir_imm_int(b, 6));
return nir_vec3(b, nir_channel(b, &tex->dest.ssa, 0),
nir_channel(b, &tex->dest.ssa, 1),
cube_array_dim);
}
static nir_ssa_def *
static nir_def *
lower_cube_image_size(nir_builder *b, nir_intrinsic_instr *intr)
{
b->cursor = nir_after_instr(&intr->instr);
if (!nir_intrinsic_image_array(intr))
return nir_trim_vector(b, &intr->dest.ssa, 2);
nir_ssa_def *array_dim = nir_channel(b, &intr->dest.ssa, 2);
nir_ssa_def *cube_array_dim = nir_idiv(b, array_dim, nir_imm_int(b, 6));
nir_def *array_dim = nir_channel(b, &intr->dest.ssa, 2);
nir_def *cube_array_dim = nir_idiv(b, array_dim, nir_imm_int(b, 6));
return nir_vec3(b, nir_channel(b, &intr->dest.ssa, 0),
nir_channel(b, &intr->dest.ssa, 1),
cube_array_dim);
@@ -482,7 +482,7 @@ make_2darray_from_cubemap_with_array(const struct glsl_type *type, bool is_image
return make_2darray_sampler_from_cubemap(type);
}
static nir_ssa_def *
static nir_def *
lower_int_cubemap_to_array_tex(nir_builder *b, nir_tex_instr *tex)
{
switch (tex->op) {
@@ -500,7 +500,7 @@ lower_int_cubemap_to_array_tex(nir_builder *b, nir_tex_instr *tex)
}
}
static nir_ssa_def *
static nir_def *
lower_cube_image_intrinsic(nir_builder *b, nir_intrinsic_instr *intr)
{
if (intr->intrinsic == nir_intrinsic_image_size ||
@@ -510,7 +510,7 @@ lower_cube_image_intrinsic(nir_builder *b, nir_intrinsic_instr *intr)
return lower_cube_image_load_store_atomic(b, intr);
}
static nir_ssa_def *
static nir_def *
lower_cube_image_deref(nir_builder *b, nir_deref_instr *deref)
{
deref->type = make_2darray_from_cubemap_with_array(
@@ -519,7 +519,7 @@ lower_cube_image_deref(nir_builder *b, nir_deref_instr *deref)
return NIR_LOWER_INSTR_PROGRESS;
}
static nir_ssa_def *
static nir_def *
lower_int_cubmap_to_array_impl(nir_builder *b, nir_instr *instr,
void *options)
{
src/microsoft/compiler/dxil_nir_lower_int_samplers.c
+43 -43
View File
@@ -49,7 +49,7 @@ lower_sample_to_txf_for_integer_tex_filter(const nir_instr *instr,
return (tex->dest_type & (nir_type_int | nir_type_uint));
}
static nir_ssa_def *
static nir_def *
dx_get_texture_lod(nir_builder *b, nir_tex_instr *tex)
{
nir_tex_instr *tql;
@@ -85,7 +85,7 @@ dx_get_texture_lod(nir_builder *b, nir_tex_instr *tex)
* check for is_array though, in the worst case we create an additional
* move the the optimization will remove later again. */
int coord_index = nir_tex_instr_src_index(tex, nir_tex_src_coord);
nir_ssa_def *ssa_src = nir_trim_vector(b, tex->src[coord_index].src.ssa,
nir_def *ssa_src = nir_trim_vector(b, tex->src[coord_index].src.ssa,
coord_components);
nir_src src = nir_src_for_ssa(ssa_src);
nir_src_copy(&tql->src[0].src, &src, &tql->instr);
@@ -113,19 +113,19 @@ dx_get_texture_lod(nir_builder *b, nir_tex_instr *tex)
}
typedef struct {
nir_ssa_def *coords;
nir_ssa_def *use_border_color;
nir_def *coords;
nir_def *use_border_color;
} wrap_result_t;
typedef struct {
nir_ssa_def *lod;
nir_ssa_def *size;
nir_def *lod;
nir_def *size;
int ncoord_comp;
wrap_result_t wrap[3];
} wrap_lower_param_t;
static void
wrap_clamp_to_edge(nir_builder *b, wrap_result_t *wrap_params, nir_ssa_def *size)
wrap_clamp_to_edge(nir_builder *b, wrap_result_t *wrap_params, nir_def *size)
{
/* clamp(coord, 0, size - 1) */
wrap_params->coords = nir_fmin(b, nir_fadd_imm(b, size, -1.0f),
@@ -133,7 +133,7 @@ wrap_clamp_to_edge(nir_builder *b, wrap_result_t *wrap_params, nir_ssa_def *size
}
static void
wrap_repeat(nir_builder *b, wrap_result_t *wrap_params, nir_ssa_def *size)
wrap_repeat(nir_builder *b, wrap_result_t *wrap_params, nir_def *size)
{
/* mod(coord, size)
* This instruction must be exact, otherwise certain sizes result in
@@ -142,8 +142,8 @@ wrap_repeat(nir_builder *b, wrap_result_t *wrap_params, nir_ssa_def *size)
nir_instr_as_alu(wrap_params->coords->parent_instr)->exact = true;
}
static nir_ssa_def *
mirror(nir_builder *b, nir_ssa_def *coord)
static nir_def *
mirror(nir_builder *b, nir_def *coord)
{
/* coord if >= 0, otherwise -(1 + coord) */
return nir_bcsel(b, nir_fge_imm(b, coord, 0.0f), coord,
@@ -151,17 +151,17 @@ mirror(nir_builder *b, nir_ssa_def *coord)
}
static void
wrap_mirror_repeat(nir_builder *b, wrap_result_t *wrap_params, nir_ssa_def *size)
wrap_mirror_repeat(nir_builder *b, wrap_result_t *wrap_params, nir_def *size)
{
/* (size − 1) − mirror(mod(coord, 2 * size) − size) */
nir_ssa_def *coord_mod2size = nir_fmod(b, wrap_params->coords, nir_fmul_imm(b, size, 2.0f));
nir_def *coord_mod2size = nir_fmod(b, wrap_params->coords, nir_fmul_imm(b, size, 2.0f));
nir_instr_as_alu(coord_mod2size->parent_instr)->exact = true;
nir_ssa_def *a = nir_fsub(b, coord_mod2size, size);
nir_def *a = nir_fsub(b, coord_mod2size, size);
wrap_params->coords = nir_fsub(b, nir_fadd_imm(b, size, -1.0f), mirror(b, a));
}
static void
wrap_mirror_clamp_to_edge(nir_builder *b, wrap_result_t *wrap_params, nir_ssa_def *size)
wrap_mirror_clamp_to_edge(nir_builder *b, wrap_result_t *wrap_params, nir_def *size)
{
/* clamp(mirror(coord), 0, size - 1) */
wrap_params->coords = nir_fmin(b, nir_fadd_imm(b, size, -1.0f),
@@ -169,19 +169,19 @@ wrap_mirror_clamp_to_edge(nir_builder *b, wrap_result_t *wrap_params, nir_ssa_de
}
static void
wrap_clamp(nir_builder *b, wrap_result_t *wrap_params, nir_ssa_def *size)
wrap_clamp(nir_builder *b, wrap_result_t *wrap_params, nir_def *size)
{
nir_ssa_def *is_low = nir_flt_imm(b, wrap_params->coords, 0.0);
nir_ssa_def *is_high = nir_fge(b, wrap_params->coords, size);
nir_def *is_low = nir_flt_imm(b, wrap_params->coords, 0.0);
nir_def *is_high = nir_fge(b, wrap_params->coords, size);
wrap_params->use_border_color = nir_ior(b, is_low, is_high);
}
static void
wrap_mirror_clamp(nir_builder *b, wrap_result_t *wrap_params, nir_ssa_def *size)
wrap_mirror_clamp(nir_builder *b, wrap_result_t *wrap_params, nir_def *size)
{
/* We have to take care of the boundaries */
nir_ssa_def *is_low = nir_flt(b, wrap_params->coords, nir_fmul_imm(b, size, -1.0));
nir_ssa_def *is_high = nir_flt(b, nir_fmul_imm(b, size, 2.0), wrap_params->coords);
nir_def *is_low = nir_flt(b, wrap_params->coords, nir_fmul_imm(b, size, -1.0));
nir_def *is_high = nir_flt(b, nir_fmul_imm(b, size, 2.0), wrap_params->coords);
wrap_params->use_border_color = nir_ior(b, is_low, is_high);
/* Within the boundaries this acts like mirror_repeat */
@@ -190,8 +190,8 @@ wrap_mirror_clamp(nir_builder *b, wrap_result_t *wrap_params, nir_ssa_def *size)
}
static wrap_result_t
wrap_coords(nir_builder *b, nir_ssa_def *coords, enum pipe_tex_wrap wrap,
nir_ssa_def *size)
wrap_coords(nir_builder *b, nir_def *coords, enum pipe_tex_wrap wrap,
nir_def *size)
{
wrap_result_t result = {coords, nir_imm_false(b)};
@@ -220,7 +220,7 @@ wrap_coords(nir_builder *b, nir_ssa_def *coords, enum pipe_tex_wrap wrap,
return result;
}
static nir_ssa_def *
static nir_def *
load_bordercolor(nir_builder *b, nir_tex_instr *tex, const dxil_wrap_sampler_state *active_state,
const dxil_texture_swizzle_state *tex_swizzle)
{
@@ -303,10 +303,10 @@ create_txf_from_tex(nir_builder *b, nir_tex_instr *tex)
return txf;
}
static nir_ssa_def *
static nir_def *
load_texel(nir_builder *b, nir_tex_instr *tex, wrap_lower_param_t *params)
{
nir_ssa_def *texcoord = NULL;
nir_def *texcoord = NULL;
/* Put coordinates back together */
switch (tex->coord_components) {
@@ -335,14 +335,14 @@ load_texel(nir_builder *b, nir_tex_instr *tex, wrap_lower_param_t *params)
typedef struct {
const dxil_wrap_sampler_state *aws;
float max_bias;
nir_ssa_def *size;
nir_def *size;
int ncoord_comp;
} lod_params;
static nir_ssa_def *
static nir_def *
evalute_active_lod(nir_builder *b, nir_tex_instr *tex, lod_params *params)
{
static nir_ssa_def *lod = NULL;
static nir_def *lod = NULL;
/* Later we use min_lod for clamping the LOD to a legal value */
float min_lod = MAX2(params->aws->min_lod, 0.0f);
@@ -357,12 +357,12 @@ evalute_active_lod(nir_builder *b, nir_tex_instr *tex, lod_params *params)
int ddy_index = nir_tex_instr_src_index(tex, nir_tex_src_ddy);
assert(ddx_index >= 0 && ddy_index >= 0);
nir_ssa_def *grad = nir_fmax(b,
nir_def *grad = nir_fmax(b,
tex->src[ddx_index].src.ssa,
tex->src[ddy_index].src.ssa);
nir_ssa_def *r = nir_fmul(b, grad, nir_i2f32(b, params->size));
nir_ssa_def *rho = nir_channel(b, r, 0);
nir_def *r = nir_fmul(b, grad, nir_i2f32(b, params->size));
nir_def *rho = nir_channel(b, r, 0);
for (int i = 1; i < params->ncoord_comp; ++i)
rho = nir_fmax(b, rho, nir_channel(b, r, i));
lod = nir_flog2(b, rho);
@@ -391,7 +391,7 @@ evalute_active_lod(nir_builder *b, nir_tex_instr *tex, lod_params *params)
* in compatibility contexts and as bias_texobj in core contexts, hence the
* implementation here is the same in both cases.
*/
nir_ssa_def *lod_bias = nir_imm_float(b, params->aws->lod_bias);
nir_def *lod_bias = nir_imm_float(b, params->aws->lod_bias);
if (unlikely(tex->op == nir_texop_txb)) {
int bias_index = nir_tex_instr_src_index(tex, nir_tex_src_bias);
@@ -418,7 +418,7 @@ evalute_active_lod(nir_builder *b, nir_tex_instr *tex, lod_params *params)
}
static nir_ssa_def *
static nir_def *
lower_sample_to_txf_for_integer_tex_impl(nir_builder *b, nir_instr *instr,
void *options)
{
@@ -443,14 +443,14 @@ lower_sample_to_txf_for_integer_tex_impl(nir_builder *b, nir_instr *instr,
b->cursor = nir_before_instr(instr);
int coord_index = nir_tex_instr_src_index(tex, nir_tex_src_coord);
nir_ssa_def *old_coord = tex->src[coord_index].src.ssa;
nir_def *old_coord = tex->src[coord_index].src.ssa;
params.ncoord_comp = tex->coord_components;
if (tex->is_array)
params.ncoord_comp -= 1;
/* This helper to get the texture size always uses LOD 0, and DirectX doesn't support
* giving another LOD when querying the texture size */
nir_ssa_def *size0 = nir_get_texture_size(b, tex);
nir_def *size0 = nir_get_texture_size(b, tex);
params.lod = nir_imm_int(b, 0);
@@ -470,16 +470,16 @@ lower_sample_to_txf_for_integer_tex_impl(nir_builder *b, nir_instr *instr,
params.size = nir_i2f32(b, size0);
}
nir_ssa_def *new_coord = old_coord;
nir_def *new_coord = old_coord;
if (!active_wrap_state->is_nonnormalized_coords) {
/* Evaluate the integer lookup coordinates for the requested LOD, don't touch the
* array index */
if (!tex->is_array) {
new_coord = nir_fmul(b, params.size, old_coord);
} else {
nir_ssa_def *array_index = nir_channel(b, old_coord, params.ncoord_comp);
nir_def *array_index = nir_channel(b, old_coord, params.ncoord_comp);
int mask = (1 << params.ncoord_comp) - 1;
nir_ssa_def *coord = nir_fmul(b, nir_channels(b, params.size, mask),
nir_def *coord = nir_fmul(b, nir_channels(b, params.size, mask),
nir_channels(b, old_coord, mask));
switch (params.ncoord_comp) {
case 1:
@@ -496,7 +496,7 @@ lower_sample_to_txf_for_integer_tex_impl(nir_builder *b, nir_instr *instr,
}
}
nir_ssa_def *coord_help[3];
nir_def *coord_help[3];
for (int i = 0; i < params.ncoord_comp; ++i)
coord_help[i] = nir_ffloor(b, nir_channel(b, new_coord, i));
@@ -507,12 +507,12 @@ lower_sample_to_txf_for_integer_tex_impl(nir_builder *b, nir_instr *instr,
/* Correct the texture coordinates for the offsets. */
int offset_index = nir_tex_instr_src_index(tex, nir_tex_src_offset);
if (offset_index >= 0) {
nir_ssa_def *offset = tex->src[offset_index].src.ssa;
nir_def *offset = tex->src[offset_index].src.ssa;
for (int i = 0; i < params.ncoord_comp; ++i)
coord_help[i] = nir_fadd(b, coord_help[i], nir_i2f32(b, nir_channel(b, offset, i)));
}
nir_ssa_def *use_border_color = nir_imm_false(b);
nir_def *use_border_color = nir_imm_false(b);
if (!active_wrap_state->skip_boundary_conditions) {
@@ -544,9 +544,9 @@ lower_sample_to_txf_for_integer_tex_impl(nir_builder *b, nir_instr *instr,
&states->tex_swizzles[tex->sampler_index]:
&one2one;
nir_ssa_def *border_color = load_bordercolor(b, tex, active_wrap_state, swizzle);
nir_def *border_color = load_bordercolor(b, tex, active_wrap_state, swizzle);
nir_if *border_else = nir_push_else(b, border_if);
nir_ssa_def *sampler_color = load_texel(b, tex, &params);
nir_def *sampler_color = load_texel(b, tex, &params);
nir_pop_if(b, border_else);
return nir_if_phi(b, border_color, sampler_color);
src/microsoft/compiler/dxil_nir_lower_vs_vertex_conversion.c
+18 -18
View File
@@ -48,25 +48,25 @@ lower_vs_vertex_conversion_filter(const nir_instr *instr, const void *options)
(get_input_target_format(var, options) != PIPE_FORMAT_NONE);
}
typedef nir_ssa_def *
(*shift_right_func)(nir_builder *build, nir_ssa_def *src0, nir_ssa_def *src1);
typedef nir_def *
(*shift_right_func)(nir_builder *build, nir_def *src0, nir_def *src1);
/* decoding the signed vs unsigned scaled format is handled
* by applying the signed or unsigned shift right function
* accordingly */
static nir_ssa_def *
from_10_10_10_2_scaled(nir_builder *b, nir_ssa_def *src,
nir_ssa_def *lshift, shift_right_func shr)
static nir_def *
from_10_10_10_2_scaled(nir_builder *b, nir_def *src,
nir_def *lshift, shift_right_func shr)
{
nir_ssa_def *rshift = nir_imm_ivec4(b, 22, 22, 22, 30);
nir_def *rshift = nir_imm_ivec4(b, 22, 22, 22, 30);
return nir_i2f32(b, shr(b, nir_ishl(b, src, lshift), rshift));
}
static nir_ssa_def *
from_10_10_10_2_snorm(nir_builder *b, nir_ssa_def *src, nir_ssa_def *lshift)
static nir_def *
from_10_10_10_2_snorm(nir_builder *b, nir_def *src, nir_def *lshift)
{
nir_ssa_def *split = from_10_10_10_2_scaled(b, src, lshift, nir_ishr);
nir_ssa_def *scale_rgb = nir_imm_vec4(b,
nir_def *split = from_10_10_10_2_scaled(b, src, lshift, nir_ishr);
nir_def *scale_rgb = nir_imm_vec4(b,
1.0f / 0x1ff,
1.0f / 0x1ff,
1.0f / 0x1ff,
@@ -74,11 +74,11 @@ from_10_10_10_2_snorm(nir_builder *b, nir_ssa_def *src, nir_ssa_def *lshift)
return nir_fmul(b, split, scale_rgb);
}
static nir_ssa_def *
from_10_10_10_2_unorm(nir_builder *b, nir_ssa_def *src, nir_ssa_def *lshift)
static nir_def *
from_10_10_10_2_unorm(nir_builder *b, nir_def *src, nir_def *lshift)
{
nir_ssa_def *split = from_10_10_10_2_scaled(b, src, lshift, nir_ushr);
nir_ssa_def *scale_rgb = nir_imm_vec4(b,
nir_def *split = from_10_10_10_2_scaled(b, src, lshift, nir_ushr);
nir_def *scale_rgb = nir_imm_vec4(b,
1.0f / 0x3ff,
1.0f / 0x3ff,
1.0f / 0x3ff,
@@ -86,19 +86,19 @@ from_10_10_10_2_unorm(nir_builder *b, nir_ssa_def *src, nir_ssa_def *lshift)
return nir_fmul(b, split, scale_rgb);
}
inline static nir_ssa_def *
inline static nir_def *
lshift_rgba(nir_builder *b)
{
return nir_imm_ivec4(b, 22, 12, 2, 0);
}
inline static nir_ssa_def *
inline static nir_def *
lshift_bgra(nir_builder *b)
{
return nir_imm_ivec4(b, 2, 12, 22, 0);
}
static nir_ssa_def *
static nir_def *
lower_vs_vertex_conversion_impl(nir_builder *b, nir_instr *instr, void *options)
{
nir_intrinsic_instr *intr = nir_instr_as_intrinsic(instr);
@@ -115,7 +115,7 @@ lower_vs_vertex_conversion_impl(nir_builder *b, nir_instr *instr, void *options)
return NULL;
return nir_vector_insert_imm(b, &intr->dest.ssa, nir_imm_int(b, 1), 3);
} else {
nir_ssa_def *src = nir_channel(b, &intr->dest.ssa, 0);
nir_def *src = nir_channel(b, &intr->dest.ssa, 0);
switch (fmt) {
case PIPE_FORMAT_R10G10B10A2_SNORM:
src/microsoft/compiler/dxil_nir_tess.c
+4 -4
View File
@@ -128,7 +128,7 @@ get_cursor_for_instr_without_cf(nir_instr *instr)
}
struct tcs_patch_loop_state {
nir_ssa_def *deref, *count;
nir_def *deref, *count;
nir_cursor begin_cursor, end_cursor, insert_cursor;
nir_loop *loop;
};
@@ -228,7 +228,7 @@ dxil_nir_split_tess_ctrl(nir_shader *nir, nir_function **patch_const_func)
continue;
nir_foreach_use_including_if_safe(src, &intr->dest.ssa) {
b.cursor = nir_before_src(src);
nir_src_rewrite_ssa(src, nir_load_invocation_id(&b));
nir_src_rewrite(src, nir_load_invocation_id(&b));
}
nir_instr_remove(instr);
}
@@ -254,7 +254,7 @@ dxil_nir_split_tess_ctrl(nir_shader *nir, nir_function **patch_const_func)
b.cursor = state.begin_cursor = get_cursor_for_instr_without_cf(instr);
start_tcs_loop(&b, &state, loop_var_deref);
}
nir_ssa_def_rewrite_uses(&intr->dest.ssa, state.count);
nir_def_rewrite_uses(&intr->dest.ssa, state.count);
break;
}
case nir_intrinsic_barrier:
@@ -313,7 +313,7 @@ remove_tess_level_accesses(nir_builder *b, nir_instr *instr, void *_data)
} else {
b->cursor = nir_after_instr(instr);
assert(nir_dest_num_components(intr->dest) == 1);
nir_ssa_def_rewrite_uses(&intr->dest.ssa, nir_ssa_undef(b, 1, intr->dest.ssa.bit_size));
nir_def_rewrite_uses(&intr->dest.ssa, nir_undef(b, 1, intr->dest.ssa.bit_size));
}
return true;
}
src/microsoft/compiler/nir_to_dxil.c
+13 -13
View File
@@ -2120,7 +2120,7 @@ bitcast_to_float(struct ntd_context *ctx, unsigned bit_size,
}
static bool
is_phi_src(nir_ssa_def *ssa)
is_phi_src(nir_def *ssa)
{
nir_foreach_use(src, ssa)
if (src->parent_instr->type == nir_instr_type_phi)
@@ -2129,7 +2129,7 @@ is_phi_src(nir_ssa_def *ssa)
}
static void
store_ssa_def(struct ntd_context *ctx, nir_ssa_def *ssa, unsigned chan,
store_ssa_def(struct ntd_context *ctx, nir_def *ssa, unsigned chan,
const struct dxil_value *value)
{
assert(ssa->index < ctx->num_defs);
@@ -2188,7 +2188,7 @@ store_alu_dest(struct ntd_context *ctx, nir_alu_instr *alu, unsigned chan,
}
static const struct dxil_value *
get_src_ssa(struct ntd_context *ctx, const nir_ssa_def *ssa, unsigned chan)
get_src_ssa(struct ntd_context *ctx, const nir_def *ssa, unsigned chan)
{
assert(ssa->index < ctx->num_defs);
assert(chan < ssa->num_components);
@@ -2298,8 +2298,8 @@ emit_shift(struct ntd_context *ctx, nir_alu_instr *alu,
dxil_module_get_int_const(&ctx->mod, shift_mask, op0_bit_size),
0);
} else {
uint64_t val = nir_ssa_scalar_as_uint(
nir_ssa_scalar_chase_alu_src(nir_get_ssa_scalar(&alu->dest.dest.ssa, 0), 1));
uint64_t val = nir_scalar_as_uint(
nir_scalar_chase_alu_src(nir_get_ssa_scalar(&alu->dest.dest.ssa, 0), 1));
op1 = dxil_module_get_int_const(&ctx->mod, val & shift_mask, op0_bit_size);
}
@@ -2906,8 +2906,8 @@ emit_alu(struct ntd_context *ctx, nir_alu_instr *alu)
case nir_op_udiv:
if (nir_src_is_const(alu->src[1].src)) {
/* It's illegal to emit a literal divide by 0 in DXIL */
nir_ssa_scalar divisor = nir_ssa_scalar_chase_alu_src(nir_get_ssa_scalar(&alu->dest.dest.ssa, 0), 1);
if (nir_ssa_scalar_as_int(divisor) == 0) {
nir_scalar divisor = nir_scalar_chase_alu_src(nir_get_ssa_scalar(&alu->dest.dest.ssa, 0), 1);
if (nir_scalar_as_int(divisor) == 0) {
store_alu_dest(ctx, alu, 0, dxil_module_get_int_const(&ctx->mod, 0, nir_dest_bit_size(alu->dest.dest)));
return true;
}
@@ -3123,7 +3123,7 @@ static bool
emit_load_global_invocation_id(struct ntd_context *ctx,
nir_intrinsic_instr *intr)
{
nir_component_mask_t comps = nir_ssa_def_components_read(&intr->dest.ssa);
nir_component_mask_t comps = nir_def_components_read(&intr->dest.ssa);
for (int i = 0; i < nir_intrinsic_dest_components(intr); i++) {
if (comps & (1 << i)) {
@@ -3145,7 +3145,7 @@ static bool
emit_load_local_invocation_id(struct ntd_context *ctx,
nir_intrinsic_instr *intr)
{
nir_component_mask_t comps = nir_ssa_def_components_read(&intr->dest.ssa);
nir_component_mask_t comps = nir_def_components_read(&intr->dest.ssa);
for (int i = 0; i < nir_intrinsic_dest_components(intr); i++) {
if (comps & (1 << i)) {
@@ -3180,7 +3180,7 @@ static bool
emit_load_local_workgroup_id(struct ntd_context *ctx,
nir_intrinsic_instr *intr)
{
nir_component_mask_t comps = nir_ssa_def_components_read(&intr->dest.ssa);
nir_component_mask_t comps = nir_def_components_read(&intr->dest.ssa);
for (int i = 0; i < nir_intrinsic_dest_components(intr); i++) {
if (comps & (1 << i)) {
@@ -4947,7 +4947,7 @@ emit_intrinsic(struct ntd_context *ctx, nir_intrinsic_instr *intr)
}
static const struct dxil_type *
dxil_type_for_const(struct ntd_context *ctx, nir_ssa_def *def)
dxil_type_for_const(struct ntd_context *ctx, nir_def *def)
{
if (BITSET_TEST(ctx->int_types, def->index) ||
!BITSET_TEST(ctx->float_types, def->index))
@@ -5634,7 +5634,7 @@ emit_tex(struct ntd_context *ctx, nir_tex_instr *instr)
}
static bool
emit_undefined(struct ntd_context *ctx, nir_ssa_undef_instr *undef)
emit_undefined(struct ntd_context *ctx, nir_undef_instr *undef)
{
for (unsigned i = 0; i < undef->def.num_components; ++i)
store_ssa_def(ctx, &undef->def, i, dxil_module_get_int32_const(&ctx->mod, 0));
@@ -5909,7 +5909,7 @@ emit_function(struct ntd_context *ctx, nir_function *func, nir_function_impl *im
if (!ctx->phis)
return false;
nir_gather_ssa_types(impl, ctx->float_types, ctx->int_types);
nir_gather_types(impl, ctx->float_types, ctx->int_types);
if (!emit_scratch(ctx, impl))
return false;