mirror of
https://gitlab.freedesktop.org/mesa/mesa.git
synced 2026-05-09 06:48:06 +02:00
pco: Replace nir_shader_lower_instructions with nir_shader_*_pass
Signed-off-by: Caius Moldovan <caius.moldovan@imgtec.com> Reviewed-by: Simon Perretta <simon.perretta@imgtec.com> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/40390>
This commit is contained in:
Caius-Moldovan-img
Marge Bot
parent
27dbe82800
commit
daeb52d38d
6 changed files with 401 additions and 531 deletions
|
|
@ -141,7 +141,8 @@ static uint8_t vectorize_filter(const nir_instr *instr, UNUSED const void *data)
|
|||
* \param[in] data User data.
|
||||
* \return True if the instruction was found.
|
||||
*/
|
||||
static bool frag_in_scalar_filter(const nir_intrinsic_instr *intr, const void *data)
|
||||
static bool frag_in_scalar_filter(const nir_intrinsic_instr *intr,
|
||||
const void *data)
|
||||
{
|
||||
nir_shader *nir = (nir_shader *)data;
|
||||
|
||||
|
|
@ -629,9 +630,7 @@ void pco_preprocess_nir(pco_ctx *ctx, nir_shader *nir)
|
|||
NULL);
|
||||
|
||||
/* Fold constant offset srcs for IO. */
|
||||
NIR_PASS(_,
|
||||
nir,
|
||||
nir_opt_constant_folding);
|
||||
NIR_PASS(_, nir, nir_opt_constant_folding);
|
||||
|
||||
NIR_PASS(_,
|
||||
nir,
|
||||
|
|
@ -1001,36 +1000,20 @@ void pco_lower_nir(pco_ctx *ctx, nir_shader *nir, pco_data *data)
|
|||
}
|
||||
}
|
||||
|
||||
static bool is_phi_with_undefs(const nir_instr *instr,
|
||||
UNUSED const void *cb_data)
|
||||
static bool
|
||||
lower_phi_with_undefs(nir_builder *b, nir_phi_instr *phi, UNUSED void *cb_data)
|
||||
{
|
||||
if (instr->type != nir_instr_type_phi)
|
||||
return false;
|
||||
|
||||
nir_phi_instr *phi = nir_instr_as_phi(instr);
|
||||
|
||||
nir_foreach_phi_src (phi_src, phi) {
|
||||
if (nir_src_is_undef(phi_src->src))
|
||||
return true;
|
||||
}
|
||||
|
||||
return false;
|
||||
}
|
||||
|
||||
static nir_def *
|
||||
lower_phi_with_undefs(nir_builder *b, nir_instr *instr, UNUSED void *cb_data)
|
||||
{
|
||||
nir_phi_instr *phi = nir_instr_as_phi(instr);
|
||||
|
||||
bool progress = false;
|
||||
nir_foreach_phi_src (phi_src, phi) {
|
||||
if (nir_src_is_undef(phi_src->src)) {
|
||||
b->cursor = nir_after_block(phi_src->pred);
|
||||
nir_src_rewrite(&phi_src->src,
|
||||
nir_imm_intN_t(b, 0, phi_src->src.ssa->bit_size));
|
||||
progress = true;
|
||||
}
|
||||
}
|
||||
|
||||
return NIR_LOWER_INSTR_PROGRESS;
|
||||
return progress;
|
||||
}
|
||||
|
||||
static bool
|
||||
|
|
@ -1094,10 +1077,10 @@ void pco_postprocess_nir(pco_ctx *ctx, nir_shader *nir, pco_data *data)
|
|||
/* Temporary: lower phi undefs to zero because at this stage we don't want to
|
||||
* lower *all* undefs to zero, but still want to avoid undefined behaviour...
|
||||
*/
|
||||
nir_shader_lower_instructions(nir,
|
||||
is_phi_with_undefs,
|
||||
lower_phi_with_undefs,
|
||||
NULL);
|
||||
nir_shader_phi_pass(nir,
|
||||
lower_phi_with_undefs,
|
||||
nir_metadata_control_flow,
|
||||
NULL);
|
||||
|
||||
NIR_PASS(_, nir, nir_convert_from_ssa, true, false);
|
||||
NIR_PASS(_, nir, nir_opt_copy_prop);
|
||||
|
|
|
|||
|
|
@ -25,17 +25,14 @@
|
|||
* \brief Lowers an I/O instruction.
|
||||
*
|
||||
* \param[in] b NIR builder.
|
||||
* \param[in] instr NIR instruction.
|
||||
* \param[in] intr NIR intrinsic instruction.
|
||||
* \param[in] cb_data User callback data.
|
||||
* \return The replacement/lowered def.
|
||||
* \return True if progress was made.
|
||||
*/
|
||||
static nir_def *lower_io(nir_builder *b, nir_instr *instr, UNUSED void *cb_data)
|
||||
static bool
|
||||
lower_io(nir_builder *b, nir_intrinsic_instr *intr, UNUSED void *cb_data)
|
||||
{
|
||||
nir_intrinsic_instr *intr = nir_instr_as_intrinsic(instr);
|
||||
b->cursor = nir_before_instr(instr);
|
||||
|
||||
ASSERTED unsigned base = nir_intrinsic_base(intr);
|
||||
assert(!base);
|
||||
b->cursor = nir_before_instr(&intr->instr);
|
||||
|
||||
nir_src *offset_src;
|
||||
switch (intr->intrinsic) {
|
||||
|
|
@ -51,41 +48,16 @@ static nir_def *lower_io(nir_builder *b, nir_instr *instr, UNUSED void *cb_data)
|
|||
break;
|
||||
|
||||
default:
|
||||
UNREACHABLE("");
|
||||
return false;
|
||||
}
|
||||
|
||||
ASSERTED unsigned base = nir_intrinsic_base(intr);
|
||||
assert(!base);
|
||||
|
||||
/* Byte offset to DWORD offset. */
|
||||
nir_src_rewrite(offset_src, nir_ushr_imm(b, offset_src->ssa, 2));
|
||||
|
||||
return NIR_LOWER_INSTR_PROGRESS;
|
||||
}
|
||||
|
||||
/**
|
||||
* \brief Filters I/O instructions that need lowering.
|
||||
*
|
||||
* \param[in] instr NIR instruction.
|
||||
* \param[in] cb_data User callback data.
|
||||
* \return True if the instruction matches the filter.
|
||||
*/
|
||||
static bool is_lowerable_io(const nir_instr *instr, UNUSED const void *cb_data)
|
||||
{
|
||||
if (instr->type != nir_instr_type_intrinsic)
|
||||
return false;
|
||||
|
||||
nir_intrinsic_instr *intr = nir_instr_as_intrinsic(instr);
|
||||
switch (intr->intrinsic) {
|
||||
case nir_intrinsic_load_push_constant:
|
||||
case nir_intrinsic_load_shared:
|
||||
case nir_intrinsic_store_shared:
|
||||
case nir_intrinsic_shared_atomic:
|
||||
case nir_intrinsic_shared_atomic_swap:
|
||||
return true;
|
||||
|
||||
default:
|
||||
break;
|
||||
}
|
||||
|
||||
return false;
|
||||
return true;
|
||||
}
|
||||
|
||||
/**
|
||||
|
|
@ -98,8 +70,10 @@ bool pco_nir_lower_io(nir_shader *shader)
|
|||
{
|
||||
bool progress = false;
|
||||
|
||||
progress |=
|
||||
nir_shader_lower_instructions(shader, is_lowerable_io, lower_io, NULL);
|
||||
progress |= nir_shader_intrinsics_pass(shader,
|
||||
lower_io,
|
||||
nir_metadata_control_flow,
|
||||
NULL);
|
||||
|
||||
return progress;
|
||||
}
|
||||
|
|
|
|||
|
|
@ -315,14 +315,14 @@ static inline bool is_processed(nir_intrinsic_instr *intr)
|
|||
return nir_alu_type_get_base_type(type) == nir_type_invalid;
|
||||
}
|
||||
|
||||
static nir_def *lower_pfo_store(nir_builder *b,
|
||||
nir_intrinsic_instr *intr,
|
||||
struct pfo_state *state)
|
||||
static bool lower_pfo_store(nir_builder *b,
|
||||
nir_intrinsic_instr *intr,
|
||||
struct pfo_state *state)
|
||||
{
|
||||
/* Skip stores we've already processed. */
|
||||
if (is_processed(intr)) {
|
||||
util_dynarray_append(&state->stores, intr);
|
||||
return NULL;
|
||||
return false;
|
||||
}
|
||||
|
||||
nir_def *input = intr->src[0].ssa;
|
||||
|
|
@ -341,8 +341,10 @@ static nir_def *lower_pfo_store(nir_builder *b,
|
|||
b->cursor = nir_before_instr(&intr->instr);
|
||||
|
||||
enum pipe_format format = state->fs->output_formats[location];
|
||||
if (format == PIPE_FORMAT_NONE)
|
||||
return NIR_LOWER_INSTR_PROGRESS_REPLACE;
|
||||
if (format == PIPE_FORMAT_NONE) {
|
||||
nir_instr_remove(&intr->instr);
|
||||
return true;
|
||||
}
|
||||
|
||||
format = to_pbe_format(b, format, &input);
|
||||
|
||||
|
|
@ -369,17 +371,18 @@ static nir_def *lower_pfo_store(nir_builder *b,
|
|||
assert(var);
|
||||
var->type = glsl_uvec_type(output->num_components);
|
||||
|
||||
return NIR_LOWER_INSTR_PROGRESS_REPLACE;
|
||||
nir_instr_remove(&intr->instr);
|
||||
return true;
|
||||
}
|
||||
|
||||
static nir_def *lower_pfo_load(nir_builder *b,
|
||||
nir_intrinsic_instr *intr,
|
||||
struct pfo_state *state)
|
||||
static bool lower_pfo_load(nir_builder *b,
|
||||
nir_intrinsic_instr *intr,
|
||||
struct pfo_state *state)
|
||||
{
|
||||
/* Skip loads we've already processed. */
|
||||
if (is_processed(intr)) {
|
||||
util_dynarray_append(&state->loads, intr);
|
||||
return NULL;
|
||||
return false;
|
||||
}
|
||||
|
||||
unsigned base = nir_intrinsic_base(intr);
|
||||
|
|
@ -400,75 +403,55 @@ static nir_def *lower_pfo_load(nir_builder *b,
|
|||
format = state->fs->output_formats[location];
|
||||
}
|
||||
|
||||
nir_def *repl;
|
||||
if (format == PIPE_FORMAT_NONE)
|
||||
return nir_undef(b, intr->def.num_components, intr->def.bit_size);
|
||||
repl = nir_undef(b, intr->def.num_components, intr->def.bit_size);
|
||||
else {
|
||||
format = to_pbe_format(b, format, NULL);
|
||||
|
||||
format = to_pbe_format(b, format, NULL);
|
||||
nir_def *packed_comps[4];
|
||||
for (unsigned c = 0; c < ARRAY_SIZE(packed_comps); ++c) {
|
||||
packed_comps[c] = nir_load_output(b,
|
||||
1,
|
||||
32,
|
||||
offset->ssa,
|
||||
.base = base,
|
||||
.component = c,
|
||||
.dest_type = nir_type_invalid | 32,
|
||||
.io_semantics = io_semantics);
|
||||
|
||||
nir_def *packed_comps[4];
|
||||
for (unsigned c = 0; c < ARRAY_SIZE(packed_comps); ++c) {
|
||||
packed_comps[c] = nir_load_output(b,
|
||||
1,
|
||||
32,
|
||||
offset->ssa,
|
||||
.base = base,
|
||||
.component = c,
|
||||
.dest_type = nir_type_invalid | 32,
|
||||
.io_semantics = io_semantics);
|
||||
nir_intrinsic_instr *load = nir_def_as_intrinsic(packed_comps[c]);
|
||||
|
||||
nir_intrinsic_instr *load =
|
||||
nir_def_as_intrinsic(packed_comps[c]);
|
||||
util_dynarray_append(&state->loads, load);
|
||||
}
|
||||
|
||||
util_dynarray_append(&state->loads, load);
|
||||
nir_alu_type dest_type = nir_intrinsic_dest_type(intr);
|
||||
repl = unpack_from_format(b,
|
||||
packed_comps,
|
||||
dest_type,
|
||||
format,
|
||||
intr->def.num_components);
|
||||
}
|
||||
|
||||
nir_alu_type dest_type = nir_intrinsic_dest_type(intr);
|
||||
return unpack_from_format(b,
|
||||
packed_comps,
|
||||
dest_type,
|
||||
format,
|
||||
intr->def.num_components);
|
||||
}
|
||||
nir_def_rewrite_uses(&intr->def, repl);
|
||||
nir_instr_remove(&intr->instr);
|
||||
|
||||
/**
|
||||
* \brief Filters PFO-related instructions.
|
||||
*
|
||||
* \param[in] instr NIR instruction.
|
||||
* \param[in] cb_data User callback data.
|
||||
* \return True if the instruction matches the filter.
|
||||
*/
|
||||
static bool is_pfo(const nir_instr *instr, UNUSED const void *cb_data)
|
||||
{
|
||||
if (instr->type != nir_instr_type_intrinsic)
|
||||
return false;
|
||||
|
||||
nir_intrinsic_instr *intr = nir_instr_as_intrinsic(instr);
|
||||
switch (intr->intrinsic) {
|
||||
case nir_intrinsic_store_output:
|
||||
case nir_intrinsic_load_output:
|
||||
case nir_intrinsic_demote:
|
||||
case nir_intrinsic_demote_if:
|
||||
return true;
|
||||
|
||||
default:
|
||||
break;
|
||||
}
|
||||
|
||||
return false;
|
||||
return true;
|
||||
}
|
||||
|
||||
/**
|
||||
* \brief Lowers a PFO-related instruction.
|
||||
*
|
||||
* \param[in] b NIR builder.
|
||||
* \param[in] instr NIR instruction.
|
||||
* \param[in] intr NIR intrinsic instruction.
|
||||
* \param[in] cb_data User callback data.
|
||||
* \return The replacement/lowered def.
|
||||
* \return True if progress was made.
|
||||
*/
|
||||
static nir_def *lower_pfo(nir_builder *b, nir_instr *instr, void *cb_data)
|
||||
static bool lower_pfo(nir_builder *b, nir_intrinsic_instr *intr, void *cb_data)
|
||||
{
|
||||
struct pfo_state *state = cb_data;
|
||||
nir_intrinsic_instr *intr = nir_instr_as_intrinsic(instr);
|
||||
|
||||
b->cursor = nir_before_instr(&intr->instr);
|
||||
|
||||
switch (intr->intrinsic) {
|
||||
case nir_intrinsic_store_output: {
|
||||
|
|
@ -480,7 +463,8 @@ static nir_def *lower_pfo(nir_builder *b, nir_instr *instr, void *cb_data)
|
|||
assert(!state->depth_feedback_src);
|
||||
state->depth_feedback_src = nir_fsat(b, intr->src[0].ssa);
|
||||
|
||||
return NIR_LOWER_INSTR_PROGRESS_REPLACE;
|
||||
nir_instr_remove(&intr->instr);
|
||||
return true;
|
||||
}
|
||||
|
||||
if (sem.location == FRAG_RESULT_SAMPLE_MASK) {
|
||||
|
|
@ -497,7 +481,8 @@ static nir_def *lower_pfo(nir_builder *b, nir_instr *instr, void *cb_data)
|
|||
val = nir_ior(b, val, cond);
|
||||
state->last_discard_store =
|
||||
nir_build_store_reg(b, val, state->discard_cond_reg);
|
||||
return NIR_LOWER_INSTR_PROGRESS_REPLACE;
|
||||
nir_instr_remove(&intr->instr);
|
||||
return true;
|
||||
}
|
||||
|
||||
UNREACHABLE("");
|
||||
|
|
@ -510,7 +495,8 @@ static nir_def *lower_pfo(nir_builder *b, nir_instr *instr, void *cb_data)
|
|||
state->has_discards = true;
|
||||
state->last_discard_store =
|
||||
nir_build_store_reg(b, nir_imm_true(b), state->discard_cond_reg);
|
||||
return NIR_LOWER_INSTR_PROGRESS_REPLACE;
|
||||
nir_instr_remove(&intr->instr);
|
||||
return true;
|
||||
|
||||
case nir_intrinsic_demote_if: {
|
||||
state->has_discards = true;
|
||||
|
|
@ -518,7 +504,8 @@ static nir_def *lower_pfo(nir_builder *b, nir_instr *instr, void *cb_data)
|
|||
val = nir_ior(b, val, intr->src[0].ssa);
|
||||
state->last_discard_store =
|
||||
nir_build_store_reg(b, val, state->discard_cond_reg);
|
||||
return NIR_LOWER_INSTR_PROGRESS_REPLACE;
|
||||
nir_instr_remove(&intr->instr);
|
||||
return true;
|
||||
}
|
||||
|
||||
default:
|
||||
|
|
@ -668,20 +655,6 @@ static bool z_replicate(nir_shader *shader, struct pfo_state *state)
|
|||
return true;
|
||||
}
|
||||
|
||||
static bool is_frag_color_out(const nir_instr *instr,
|
||||
UNUSED const void *cb_data)
|
||||
{
|
||||
if (instr->type != nir_instr_type_intrinsic)
|
||||
return false;
|
||||
|
||||
nir_intrinsic_instr *intr = nir_instr_as_intrinsic(instr);
|
||||
if (intr->intrinsic != nir_intrinsic_store_output)
|
||||
return false;
|
||||
|
||||
gl_frag_result location = nir_intrinsic_io_semantics(intr).location;
|
||||
return location >= FRAG_RESULT_DATA0 && location < FRAG_RESULT_MAX;
|
||||
}
|
||||
|
||||
static bool lower_demote_samples(nir_builder *b,
|
||||
nir_intrinsic_instr *intr,
|
||||
UNUSED void *cb_data)
|
||||
|
|
@ -725,10 +698,16 @@ bool pco_nir_lower_alpha_to_coverage(nir_shader *shader)
|
|||
return true;
|
||||
}
|
||||
|
||||
static nir_def *
|
||||
lower_alpha_to_one(nir_builder *b, nir_instr *instr, UNUSED void *cb_data)
|
||||
static bool lower_alpha_to_one(nir_builder *b,
|
||||
nir_intrinsic_instr *intr,
|
||||
UNUSED void *cb_data)
|
||||
{
|
||||
nir_intrinsic_instr *intr = nir_instr_as_intrinsic(instr);
|
||||
if (intr->intrinsic != nir_intrinsic_store_output)
|
||||
return false;
|
||||
|
||||
gl_frag_result location = nir_intrinsic_io_semantics(intr).location;
|
||||
if (!(location >= FRAG_RESULT_DATA0 && location < FRAG_RESULT_MAX))
|
||||
return false;
|
||||
|
||||
nir_src *input_src = &intr->src[0];
|
||||
nir_def *input = input_src->ssa;
|
||||
|
|
@ -737,7 +716,7 @@ lower_alpha_to_one(nir_builder *b, nir_instr *instr, UNUSED void *cb_data)
|
|||
|
||||
/* Skip color write that don't include alpha. */
|
||||
if (input->num_components != 4)
|
||||
return NULL;
|
||||
return false;
|
||||
|
||||
b->cursor = nir_before_instr(&intr->instr);
|
||||
|
||||
|
|
@ -754,7 +733,7 @@ lower_alpha_to_one(nir_builder *b, nir_instr *instr, UNUSED void *cb_data)
|
|||
|
||||
nir_src_rewrite(input_src, nir_vector_insert_imm(b, input, alpha, 3));
|
||||
|
||||
return NIR_LOWER_INSTR_PROGRESS;
|
||||
return true;
|
||||
}
|
||||
|
||||
static bool is_load_sample_mask(const nir_instr *instr,
|
||||
|
|
@ -767,23 +746,33 @@ static bool is_load_sample_mask(const nir_instr *instr,
|
|||
return intr->intrinsic == nir_intrinsic_load_sample_mask_in;
|
||||
}
|
||||
|
||||
static nir_def *
|
||||
lower_load_sample_mask(nir_builder *b, nir_instr *instr, UNUSED void *cb_data)
|
||||
static bool lower_load_sample_mask(nir_builder *b,
|
||||
nir_intrinsic_instr *intr,
|
||||
UNUSED void *cb_data)
|
||||
{
|
||||
b->cursor = nir_before_instr(instr);
|
||||
if (intr->intrinsic != nir_intrinsic_load_sample_mask_in)
|
||||
return false;
|
||||
|
||||
b->cursor = nir_before_instr(&intr->instr);
|
||||
|
||||
nir_def *smp_msk =
|
||||
nir_ubitfield_extract_imm(b, nir_load_fs_meta_pco(b), 9, 16);
|
||||
|
||||
smp_msk = nir_iand(b, smp_msk, nir_load_savmsk_vm_pco(b));
|
||||
|
||||
return smp_msk;
|
||||
nir_def_rewrite_uses(&intr->def, smp_msk);
|
||||
nir_instr_remove(&intr->instr);
|
||||
return true;
|
||||
}
|
||||
|
||||
static nir_def *
|
||||
lower_color_write_enable(nir_builder *b, nir_instr *instr, UNUSED void *cb_data)
|
||||
static bool lower_color_write_enable(nir_builder *b,
|
||||
nir_intrinsic_instr *intr,
|
||||
UNUSED void *cb_data)
|
||||
{
|
||||
nir_intrinsic_instr *intr = nir_instr_as_intrinsic(instr);
|
||||
if (intr->intrinsic != nir_intrinsic_store_output)
|
||||
return false;
|
||||
|
||||
gl_frag_result location = nir_intrinsic_io_semantics(intr).location;
|
||||
if (!(location >= FRAG_RESULT_DATA0 && location < FRAG_RESULT_MAX))
|
||||
return false;
|
||||
|
||||
nir_src *input_src = &intr->src[0];
|
||||
nir_def *input = input_src->ssa;
|
||||
|
|
@ -819,7 +808,7 @@ lower_color_write_enable(nir_builder *b, nir_instr *instr, UNUSED void *cb_data)
|
|||
nir_src_rewrite(input_src,
|
||||
nir_bcsel(b, color_write_enabled, input, prev_input));
|
||||
|
||||
return NIR_LOWER_INSTR_PROGRESS;
|
||||
return true;
|
||||
}
|
||||
|
||||
/**
|
||||
|
|
@ -853,28 +842,32 @@ bool pco_nir_pfo(nir_shader *shader, pco_fs_data *fs)
|
|||
* the stores
|
||||
*/
|
||||
if (!shader->info.internal) {
|
||||
progress |= nir_shader_lower_instructions(shader,
|
||||
is_frag_color_out,
|
||||
lower_alpha_to_one,
|
||||
&state);
|
||||
progress |= nir_shader_intrinsics_pass(shader,
|
||||
lower_alpha_to_one,
|
||||
nir_metadata_control_flow,
|
||||
&state);
|
||||
}
|
||||
|
||||
if (fs->meta_present.color_write_enable)
|
||||
progress |= nir_shader_lower_instructions(shader,
|
||||
is_frag_color_out,
|
||||
lower_color_write_enable,
|
||||
NULL);
|
||||
if (fs->meta_present.color_write_enable) {
|
||||
progress |= nir_shader_intrinsics_pass(shader,
|
||||
lower_color_write_enable,
|
||||
nir_metadata_control_flow,
|
||||
NULL);
|
||||
}
|
||||
|
||||
progress |= nir_shader_lower_instructions(shader, is_pfo, lower_pfo, &state);
|
||||
progress |= nir_shader_intrinsics_pass(shader,
|
||||
lower_pfo,
|
||||
nir_metadata_control_flow,
|
||||
&state);
|
||||
progress |= lower_isp_fb(&b, &state);
|
||||
|
||||
progress |= sink_outputs(shader, &state);
|
||||
progress |= z_replicate(shader, &state);
|
||||
|
||||
progress |= nir_shader_lower_instructions(shader,
|
||||
is_load_sample_mask,
|
||||
lower_load_sample_mask,
|
||||
NULL);
|
||||
progress |= nir_shader_intrinsics_pass(shader,
|
||||
lower_load_sample_mask,
|
||||
nir_metadata_control_flow,
|
||||
NULL);
|
||||
|
||||
util_dynarray_fini(&state.stores);
|
||||
util_dynarray_fini(&state.loads);
|
||||
|
|
@ -882,42 +875,33 @@ bool pco_nir_pfo(nir_shader *shader, pco_fs_data *fs)
|
|||
return progress;
|
||||
}
|
||||
|
||||
static nir_def *lower_pvi(nir_builder *b, nir_instr *instr, void *cb_data)
|
||||
static bool lower_pvi(nir_builder *b, nir_intrinsic_instr *intr, void *cb_data)
|
||||
{
|
||||
nir_intrinsic_instr *intr = nir_instr_as_intrinsic(instr);
|
||||
struct pvi_state *state = cb_data;
|
||||
|
||||
unsigned start_comp = nir_intrinsic_component(intr);
|
||||
unsigned num_comps = intr->def.num_components;
|
||||
|
||||
ASSERTED nir_src *offset = &intr->src[0];
|
||||
assert(nir_src_as_uint(*offset) == 0);
|
||||
|
||||
struct nir_io_semantics io_semantics = nir_intrinsic_io_semantics(intr);
|
||||
gl_vert_attrib location = io_semantics.location;
|
||||
nir_def *attrib = state->attribs[location - VERT_ATTRIB_GENERIC0];
|
||||
assert(attrib);
|
||||
|
||||
b->cursor = nir_before_instr(&intr->instr);
|
||||
return nir_channels(b, attrib, BITFIELD_RANGE(start_comp, num_comps));
|
||||
}
|
||||
|
||||
static bool is_pvi(const nir_instr *instr, const void *cb_data)
|
||||
{
|
||||
if (instr->type != nir_instr_type_intrinsic)
|
||||
return false;
|
||||
|
||||
nir_intrinsic_instr *intr = nir_instr_as_intrinsic(instr);
|
||||
if (intr->intrinsic != nir_intrinsic_load_input)
|
||||
return false;
|
||||
|
||||
if (is_processed(intr))
|
||||
return false;
|
||||
|
||||
unsigned start_comp = nir_intrinsic_component(intr);
|
||||
unsigned num_comps = intr->def.num_components;
|
||||
|
||||
ASSERTED gl_vert_attrib location = nir_intrinsic_io_semantics(intr).location;
|
||||
assert(location >= VERT_ATTRIB_GENERIC0 &&
|
||||
location <= VERT_ATTRIB_GENERIC15);
|
||||
|
||||
ASSERTED nir_src *offset = &intr->src[0];
|
||||
assert(nir_src_as_uint(*offset) == 0);
|
||||
|
||||
nir_def *attrib = state->attribs[location - VERT_ATTRIB_GENERIC0];
|
||||
assert(attrib);
|
||||
|
||||
b->cursor = nir_before_instr(&intr->instr);
|
||||
attrib = nir_channels(b, attrib, BITFIELD_RANGE(start_comp, num_comps));
|
||||
nir_def_rewrite_uses(&intr->def, attrib);
|
||||
nir_instr_remove(&intr->instr);
|
||||
return true;
|
||||
}
|
||||
|
||||
|
|
@ -999,7 +983,10 @@ bool pco_nir_pvi(nir_shader *shader, pco_vs_data *vs)
|
|||
unpack_from_format(&b, packed_comps, base_type, format, 4);
|
||||
}
|
||||
|
||||
nir_shader_lower_instructions(shader, is_pvi, lower_pvi, &state);
|
||||
nir_shader_intrinsics_pass(shader,
|
||||
lower_pvi,
|
||||
nir_metadata_control_flow,
|
||||
&state);
|
||||
|
||||
return true;
|
||||
}
|
||||
|
|
@ -1074,24 +1061,7 @@ bool pco_nir_point_size(nir_shader *shader)
|
|||
return true;
|
||||
}
|
||||
|
||||
static bool is_fs_intr(const nir_instr *instr, UNUSED const void *cb_data)
|
||||
{
|
||||
if (instr->type != nir_instr_type_intrinsic)
|
||||
return false;
|
||||
|
||||
nir_intrinsic_instr *intr = nir_instr_as_intrinsic(instr);
|
||||
switch (intr->intrinsic) {
|
||||
case nir_intrinsic_load_front_face:
|
||||
return true;
|
||||
|
||||
default:
|
||||
break;
|
||||
}
|
||||
|
||||
return false;
|
||||
}
|
||||
|
||||
static nir_def *lower_front_face(nir_builder *b, nir_intrinsic_instr *intr)
|
||||
static bool lower_front_face(nir_builder *b, nir_intrinsic_instr *intr)
|
||||
{
|
||||
nir_def *face_ccw = nir_load_face_ccw_pco(b);
|
||||
nir_def *front_face = nir_ieq_imm(b, face_ccw, 0);
|
||||
|
|
@ -1103,32 +1073,35 @@ static nir_def *lower_front_face(nir_builder *b, nir_intrinsic_instr *intr)
|
|||
[PCO_FRONT_FACE_OP_TRUE] = nir_imm_true(b),
|
||||
};
|
||||
|
||||
return nir_select_from_ssa_def_array(b,
|
||||
ff_elems,
|
||||
ARRAY_SIZE(ff_elems),
|
||||
ff_op);
|
||||
nir_def *ff_selected =
|
||||
nir_select_from_ssa_def_array(b, ff_elems, ARRAY_SIZE(ff_elems), ff_op);
|
||||
nir_def_rewrite_uses(&intr->def, ff_selected);
|
||||
nir_instr_remove(&intr->instr);
|
||||
return true;
|
||||
}
|
||||
|
||||
static nir_def *
|
||||
lower_fs_intr(nir_builder *b, nir_instr *instr, UNUSED void *cb_data)
|
||||
static bool
|
||||
lower_fs_intr(nir_builder *b, nir_intrinsic_instr *intr, UNUSED void *cb_data)
|
||||
{
|
||||
nir_intrinsic_instr *intr = nir_instr_as_intrinsic(instr);
|
||||
switch (intr->intrinsic) {
|
||||
case nir_intrinsic_load_front_face:
|
||||
b->cursor = nir_before_instr(&intr->instr);
|
||||
return lower_front_face(b, intr);
|
||||
|
||||
default:
|
||||
break;
|
||||
}
|
||||
|
||||
UNREACHABLE("");
|
||||
return false;
|
||||
}
|
||||
|
||||
bool pco_nir_lower_fs_intrinsics(nir_shader *shader)
|
||||
{
|
||||
assert(shader->info.stage == MESA_SHADER_FRAGMENT);
|
||||
|
||||
return nir_shader_lower_instructions(shader, is_fs_intr, lower_fs_intr, NULL);
|
||||
return nir_shader_intrinsics_pass(shader,
|
||||
lower_fs_intr,
|
||||
nir_metadata_control_flow,
|
||||
NULL);
|
||||
}
|
||||
|
||||
static bool
|
||||
|
|
@ -1232,13 +1205,9 @@ clone_clip_cull_stores(nir_builder *b, nir_intrinsic_instr *intr, void *data)
|
|||
return true;
|
||||
}
|
||||
|
||||
static bool is_clip_cull_load(const nir_instr *instr,
|
||||
UNUSED const void *cb_data)
|
||||
static bool
|
||||
swap_clip_cull_load(nir_builder *b, nir_intrinsic_instr *intr, void *cb_data)
|
||||
{
|
||||
if (instr->type != nir_instr_type_intrinsic)
|
||||
return false;
|
||||
|
||||
nir_intrinsic_instr *intr = nir_instr_as_intrinsic(instr);
|
||||
if (intr->intrinsic != nir_intrinsic_load_deref)
|
||||
return false;
|
||||
|
||||
|
|
@ -1248,16 +1217,11 @@ static bool is_clip_cull_load(const nir_instr *instr,
|
|||
|
||||
nir_variable *var = nir_deref_instr_get_variable(deref);
|
||||
|
||||
return var->data.location == VARYING_SLOT_CLIP_DIST0 ||
|
||||
var->data.location == VARYING_SLOT_CLIP_DIST1;
|
||||
}
|
||||
if (var->data.location != VARYING_SLOT_CLIP_DIST0 &&
|
||||
var->data.location != VARYING_SLOT_CLIP_DIST1)
|
||||
return false;
|
||||
|
||||
static nir_def *
|
||||
swap_clip_cull_load(nir_builder *b, nir_instr *instr, void *cb_data)
|
||||
{
|
||||
nir_intrinsic_instr *intr = nir_instr_as_intrinsic(instr);
|
||||
nir_deref_instr *deref = nir_src_as_deref(intr->src[0]);
|
||||
nir_variable *var = nir_deref_instr_get_variable(deref);
|
||||
b->cursor = nir_before_instr(&intr->instr);
|
||||
|
||||
unsigned var_index = var->data.location - VARYING_SLOT_CLIP_DIST0;
|
||||
nir_def *index =
|
||||
|
|
@ -1265,7 +1229,10 @@ swap_clip_cull_load(nir_builder *b, nir_instr *instr, void *cb_data)
|
|||
index = nir_iadd_imm(b, index, var_index * 4);
|
||||
|
||||
nir_variable *clone_var = cb_data;
|
||||
return nir_load_array_var(b, clone_var, index);
|
||||
nir_def *loaded_clone_var = nir_load_array_var(b, clone_var, index);
|
||||
nir_def_rewrite_uses(&intr->def, loaded_clone_var);
|
||||
nir_instr_remove(&intr->instr);
|
||||
return true;
|
||||
}
|
||||
|
||||
bool pco_nir_link_clip_cull_vars(nir_shader *producer, nir_shader *consumer)
|
||||
|
|
@ -1304,10 +1271,10 @@ bool pco_nir_link_clip_cull_vars(nir_shader *producer, nir_shader *consumer)
|
|||
nir_variable_create(consumer, nir_var_shader_in, clone_var_type, NULL);
|
||||
clone_var->data.location = clone_slot;
|
||||
|
||||
nir_shader_lower_instructions(consumer,
|
||||
is_clip_cull_load,
|
||||
swap_clip_cull_load,
|
||||
clone_var);
|
||||
nir_shader_intrinsics_pass(consumer,
|
||||
swap_clip_cull_load,
|
||||
nir_metadata_control_flow,
|
||||
clone_var);
|
||||
|
||||
return true;
|
||||
}
|
||||
|
|
|
|||
|
|
@ -28,24 +28,29 @@
|
|||
* \brief Lowers a barrier instruction.
|
||||
*
|
||||
* \param[in] b NIR builder.
|
||||
* \param[in] instr NIR instruction.
|
||||
* \param[in] intr NIR intrinsic instruction.
|
||||
* \param[in] cb_data User callback data.
|
||||
* \return The replacement/lowered def.
|
||||
* \return True if progress was made.
|
||||
*/
|
||||
static nir_def *lower_barrier(nir_builder *b, nir_instr *instr, void *cb_data)
|
||||
static bool
|
||||
lower_barrier(nir_builder *b, nir_intrinsic_instr *intr, void *cb_data)
|
||||
{
|
||||
if (intr->intrinsic != nir_intrinsic_barrier)
|
||||
return false;
|
||||
|
||||
struct shader_info *info = &b->shader->info;
|
||||
bool *uses_usclib = cb_data;
|
||||
|
||||
nir_intrinsic_instr *intr = nir_instr_as_intrinsic(instr);
|
||||
mesa_scope exec_scope = nir_intrinsic_execution_scope(intr);
|
||||
|
||||
unsigned wg_size = info->workgroup_size[0] * info->workgroup_size[1] *
|
||||
info->workgroup_size[2];
|
||||
|
||||
if (wg_size <= ROGUE_MAX_INSTANCES_PER_TASK || exec_scope == SCOPE_NONE ||
|
||||
exec_scope == SCOPE_SUBGROUP)
|
||||
return NIR_LOWER_INSTR_PROGRESS_REPLACE;
|
||||
exec_scope == SCOPE_SUBGROUP) {
|
||||
nir_instr_remove(&intr->instr);
|
||||
return true;
|
||||
}
|
||||
|
||||
/* TODO: We might be able to re-use barrier counters. */
|
||||
unsigned counter_offset = info->shared_size;
|
||||
|
|
@ -56,25 +61,10 @@ static nir_def *lower_barrier(nir_builder *b, nir_instr *instr, void *cb_data)
|
|||
|
||||
unsigned num_slots = DIV_ROUND_UP(wg_size, ROGUE_MAX_INSTANCES_PER_TASK);
|
||||
|
||||
b->cursor = nir_before_instr(instr);
|
||||
b->cursor = nir_before_instr(&intr->instr);
|
||||
usclib_barrier(b, nir_imm_int(b, num_slots), nir_imm_int(b, counter_offset));
|
||||
|
||||
return NIR_LOWER_INSTR_PROGRESS_REPLACE;
|
||||
}
|
||||
|
||||
/**
|
||||
* \brief Filters barrier instructions.
|
||||
*
|
||||
* \param[in] instr NIR instruction.
|
||||
* \param[in] cb_data User callback data.
|
||||
* \return True if the instruction matches the filter.
|
||||
*/
|
||||
static bool is_barrier(const nir_instr *instr, UNUSED const void *cb_data)
|
||||
{
|
||||
if (instr->type != nir_instr_type_intrinsic)
|
||||
return false;
|
||||
|
||||
return nir_instr_as_intrinsic(instr)->intrinsic == nir_intrinsic_barrier;
|
||||
nir_instr_remove(&intr->instr);
|
||||
return true;
|
||||
}
|
||||
|
||||
/**
|
||||
|
|
@ -85,23 +75,26 @@ static bool is_barrier(const nir_instr *instr, UNUSED const void *cb_data)
|
|||
*/
|
||||
bool pco_nir_lower_barriers(nir_shader *shader, pco_data *data)
|
||||
{
|
||||
bool progress = nir_shader_lower_instructions(shader,
|
||||
is_barrier,
|
||||
lower_barrier,
|
||||
&data->common.uses.usclib);
|
||||
bool progress = nir_shader_intrinsics_pass(shader,
|
||||
lower_barrier,
|
||||
nir_metadata_none,
|
||||
&data->common.uses.usclib);
|
||||
|
||||
data->common.uses.barriers |= progress;
|
||||
|
||||
return progress;
|
||||
}
|
||||
|
||||
static nir_def *
|
||||
lower_usclib_atomic(nir_builder *b, nir_instr *instr, void *cb_data)
|
||||
static bool
|
||||
lower_usclib_atomic(nir_builder *b, nir_intrinsic_instr *intr, void *cb_data)
|
||||
{
|
||||
nir_intrinsic_instr *intr = nir_instr_as_intrinsic(instr);
|
||||
if (intr->intrinsic != nir_intrinsic_ssbo_atomic_swap &&
|
||||
intr->intrinsic != nir_intrinsic_global_atomic_swap_pco)
|
||||
return false;
|
||||
|
||||
bool *uses_usclib = cb_data;
|
||||
|
||||
b->cursor = nir_before_instr(instr);
|
||||
b->cursor = nir_before_instr(&intr->instr);
|
||||
|
||||
if (intr->intrinsic == nir_intrinsic_ssbo_atomic_swap) {
|
||||
nir_def *buffer = intr->src[0].ssa;
|
||||
|
|
@ -116,11 +109,11 @@ lower_usclib_atomic(nir_builder *b, nir_instr *instr, void *cb_data)
|
|||
assert(num_components == 1 && bit_size == 32);
|
||||
|
||||
*uses_usclib = true;
|
||||
return usclib_emu_ssbo_atomic_comp_swap(b,
|
||||
buffer,
|
||||
offset,
|
||||
value,
|
||||
value_swap);
|
||||
nir_def *emulated =
|
||||
usclib_emu_ssbo_atomic_comp_swap(b, buffer, offset, value, value_swap);
|
||||
nir_def_rewrite_uses(&intr->def, emulated);
|
||||
nir_instr_remove(&intr->instr);
|
||||
return true;
|
||||
}
|
||||
|
||||
nir_def *addr_data = intr->src[0].ssa;
|
||||
|
|
@ -133,7 +126,12 @@ lower_usclib_atomic(nir_builder *b, nir_instr *instr, void *cb_data)
|
|||
assert(num_components == 1 && bit_size == 32);
|
||||
|
||||
*uses_usclib = true;
|
||||
return usclib_emu_global_atomic_comp_swap(b, addr, value, value_swap);
|
||||
|
||||
nir_def *emulated =
|
||||
usclib_emu_global_atomic_comp_swap(b, addr, value, value_swap);
|
||||
nir_def_rewrite_uses(&intr->def, emulated);
|
||||
nir_instr_remove(&intr->instr);
|
||||
return true;
|
||||
}
|
||||
|
||||
static bool lower_global_atomic_intrinsic(nir_builder *b,
|
||||
|
|
@ -184,25 +182,6 @@ static bool lower_global_atomic_intrinsic(nir_builder *b,
|
|||
return true;
|
||||
}
|
||||
|
||||
/**
|
||||
* \brief Filters atomic instructions emulated with usclib.
|
||||
*
|
||||
* \param[in] instr NIR instruction.
|
||||
* \param[in] cb_data User callback data.
|
||||
* \return True if the instruction matches the filter.
|
||||
*/
|
||||
static bool atomic_uses_usclib(const nir_instr *instr,
|
||||
UNUSED const void *cb_data)
|
||||
{
|
||||
if (instr->type != nir_instr_type_intrinsic)
|
||||
return false;
|
||||
|
||||
nir_intrinsic_instr *intr = nir_instr_as_intrinsic(instr);
|
||||
|
||||
return intr->intrinsic == nir_intrinsic_ssbo_atomic_swap ||
|
||||
intr->intrinsic == nir_intrinsic_global_atomic_swap_pco;
|
||||
}
|
||||
|
||||
/**
|
||||
* \brief Atomics lowering pass.
|
||||
*
|
||||
|
|
@ -217,70 +196,58 @@ bool pco_nir_lower_atomics(nir_shader *shader, pco_data *data)
|
|||
lower_global_atomic_intrinsic,
|
||||
nir_metadata_none,
|
||||
NULL);
|
||||
progress |= nir_shader_lower_instructions(shader,
|
||||
atomic_uses_usclib,
|
||||
lower_usclib_atomic,
|
||||
&data->common.uses.usclib);
|
||||
progress |= nir_shader_intrinsics_pass(shader,
|
||||
lower_usclib_atomic,
|
||||
nir_metadata_none,
|
||||
&data->common.uses.usclib);
|
||||
|
||||
return progress;
|
||||
}
|
||||
|
||||
static nir_def *
|
||||
lower_subgroup_intrinsic(nir_builder *b, nir_instr *instr, void *cb_data)
|
||||
static bool lower_subgroup_intrinsic(nir_builder *b,
|
||||
nir_intrinsic_instr *intr,
|
||||
void *cb_data)
|
||||
{
|
||||
nir_intrinsic_instr *intr = nir_instr_as_intrinsic(instr);
|
||||
assert(intr->def.num_components == 1);
|
||||
nir_def *new_def;
|
||||
|
||||
b->cursor = nir_before_instr(&intr->instr);
|
||||
|
||||
switch (intr->intrinsic) {
|
||||
case nir_intrinsic_load_subgroup_size:
|
||||
return nir_imm_int(b, 1);
|
||||
new_def = nir_imm_int(b, 1);
|
||||
break;
|
||||
|
||||
case nir_intrinsic_load_subgroup_invocation:
|
||||
return nir_imm_int(b, 0);
|
||||
new_def = nir_imm_int(b, 0);
|
||||
break;
|
||||
|
||||
case nir_intrinsic_load_num_subgroups:
|
||||
return nir_imm_int(b,
|
||||
b->shader->info.workgroup_size[0] *
|
||||
b->shader->info.workgroup_size[1] *
|
||||
b->shader->info.workgroup_size[2]);
|
||||
new_def = nir_imm_int(b,
|
||||
b->shader->info.workgroup_size[0] *
|
||||
b->shader->info.workgroup_size[1] *
|
||||
b->shader->info.workgroup_size[2]);
|
||||
break;
|
||||
|
||||
case nir_intrinsic_load_subgroup_id:
|
||||
return nir_load_local_invocation_index(b);
|
||||
new_def = nir_load_local_invocation_index(b);
|
||||
break;
|
||||
|
||||
case nir_intrinsic_first_invocation:
|
||||
return nir_imm_int(b, 0);
|
||||
new_def = nir_imm_int(b, 0);
|
||||
break;
|
||||
|
||||
case nir_intrinsic_elect:
|
||||
return nir_imm_true(b);
|
||||
new_def = nir_imm_true(b);
|
||||
break;
|
||||
|
||||
default:
|
||||
break;
|
||||
}
|
||||
|
||||
UNREACHABLE("");
|
||||
}
|
||||
|
||||
static bool is_subgroup_intrinsic(const nir_instr *instr,
|
||||
UNUSED const void *cb_data)
|
||||
{
|
||||
if (instr->type != nir_instr_type_intrinsic)
|
||||
return false;
|
||||
|
||||
nir_intrinsic_instr *intr = nir_instr_as_intrinsic(instr);
|
||||
switch (intr->intrinsic) {
|
||||
case nir_intrinsic_load_subgroup_size:
|
||||
case nir_intrinsic_load_subgroup_invocation:
|
||||
case nir_intrinsic_load_num_subgroups:
|
||||
case nir_intrinsic_load_subgroup_id:
|
||||
case nir_intrinsic_first_invocation:
|
||||
case nir_intrinsic_elect:
|
||||
return true;
|
||||
|
||||
default:
|
||||
break;
|
||||
}
|
||||
|
||||
return false;
|
||||
nir_def_rewrite_uses(&intr->def, new_def);
|
||||
nir_instr_remove(&intr->instr);
|
||||
assert(intr->def.num_components == 1);
|
||||
return true;
|
||||
}
|
||||
|
||||
bool pco_nir_lower_subgroups(nir_shader *shader)
|
||||
|
|
@ -289,8 +256,8 @@ bool pco_nir_lower_subgroups(nir_shader *shader)
|
|||
shader->info.min_subgroup_size = 1;
|
||||
shader->info.max_subgroup_size = 1;
|
||||
|
||||
return nir_shader_lower_instructions(shader,
|
||||
is_subgroup_intrinsic,
|
||||
lower_subgroup_intrinsic,
|
||||
NULL);
|
||||
return nir_shader_intrinsics_pass(shader,
|
||||
lower_subgroup_intrinsic,
|
||||
nir_metadata_control_flow,
|
||||
NULL);
|
||||
}
|
||||
|
|
|
|||
|
|
@ -44,29 +44,36 @@ static inline nir_def *get_src_def(nir_tex_instr *tex,
|
|||
* \param[in] b NIR builder.
|
||||
* \param[in] tex NIR texture instruction.
|
||||
* \param[in] tex_state Texture state words.
|
||||
* \return The replacement/lowered def.
|
||||
* \return True if progress was made.
|
||||
*/
|
||||
static nir_def *lower_tex_query_basic(nir_builder *b,
|
||||
nir_tex_instr *tex,
|
||||
nir_def *tex_state,
|
||||
nir_def *tex_meta,
|
||||
pco_data *data)
|
||||
static bool lower_tex_query_basic(nir_builder *b,
|
||||
nir_tex_instr *tex,
|
||||
nir_def *tex_state,
|
||||
nir_def *tex_meta,
|
||||
pco_data *data)
|
||||
{
|
||||
nir_def *new_def;
|
||||
|
||||
b->cursor = nir_before_instr(&tex->instr);
|
||||
|
||||
switch (tex->op) {
|
||||
case nir_texop_query_levels:
|
||||
data->common.uses.usclib = true;
|
||||
return usclib_tex_state_levels(b, tex_state);
|
||||
new_def = usclib_tex_state_levels(b, tex_state);
|
||||
break;
|
||||
|
||||
case nir_texop_texture_samples:
|
||||
data->common.uses.usclib = true;
|
||||
return usclib_tex_state_samples(b, tex_state);
|
||||
new_def = usclib_tex_state_samples(b, tex_state);
|
||||
break;
|
||||
|
||||
case nir_texop_txs: {
|
||||
if (tex->sampler_dim == GLSL_SAMPLER_DIM_BUF) {
|
||||
assert(tex->def.num_components == 1);
|
||||
assert(!tex->is_array);
|
||||
|
||||
return nir_channel(b, tex_meta, PCO_IMAGE_META_BUFFER_ELEMS);
|
||||
new_def = nir_channel(b, tex_meta, PCO_IMAGE_META_BUFFER_ELEMS);
|
||||
break;
|
||||
}
|
||||
|
||||
nir_def *num_comps = nir_imm_int(b, tex->def.num_components);
|
||||
|
|
@ -84,14 +91,16 @@ static nir_def *lower_tex_query_basic(nir_builder *b,
|
|||
|
||||
data->common.uses.usclib = true;
|
||||
|
||||
return nir_trim_vector(b, size_comps, tex->def.num_components);
|
||||
}
|
||||
|
||||
default:
|
||||
new_def = nir_trim_vector(b, size_comps, tex->def.num_components);
|
||||
break;
|
||||
}
|
||||
|
||||
UNREACHABLE("");
|
||||
default:
|
||||
UNREACHABLE("");
|
||||
}
|
||||
|
||||
nir_def_rewrite_uses(&tex->def, new_def);
|
||||
return true;
|
||||
}
|
||||
|
||||
static inline enum pco_dim to_pco_dim(enum glsl_sampler_dim dim)
|
||||
|
|
@ -442,13 +451,12 @@ static nir_def *lower_tex_shadow(nir_builder *b,
|
|||
* \brief Lowers a texture instruction.
|
||||
*
|
||||
* \param[in] b NIR builder.
|
||||
* \param[in] instr NIR instruction.
|
||||
* \param[in] instr NIR texture instruction.
|
||||
* \param[in] cb_data User callback data.
|
||||
* \return The replacement/lowered def.
|
||||
* \return True if progress was made.
|
||||
*/
|
||||
static nir_def *lower_tex(nir_builder *b, nir_instr *instr, void *cb_data)
|
||||
static bool lower_tex(nir_builder *b, nir_tex_instr *tex, void *cb_data)
|
||||
{
|
||||
nir_tex_instr *tex = nir_instr_as_tex(instr);
|
||||
struct state *state = cb_data;
|
||||
pco_data *data = state->data;
|
||||
pco_ctx *ctx = state->ctx;
|
||||
|
|
@ -468,7 +476,7 @@ static nir_def *lower_tex(nir_builder *b, nir_instr *instr, void *cb_data)
|
|||
? PVR_HAS_FEATURE(dev_info, tpu_extended_integer_lookup)
|
||||
: false;
|
||||
|
||||
b->cursor = nir_before_instr(instr);
|
||||
b->cursor = nir_before_instr(&tex->instr);
|
||||
|
||||
/* Process tex sources, build up the smp flags and data words. */
|
||||
BITSET_DECLARE(tex_src_set, nir_num_tex_src_types) = { 0 };
|
||||
|
|
@ -771,20 +779,8 @@ static nir_def *lower_tex(nir_builder *b, nir_instr *instr, void *cb_data)
|
|||
|
||||
result = lower_tex_shadow(b, result, comparator, compare_op);
|
||||
}
|
||||
|
||||
return result;
|
||||
}
|
||||
|
||||
/**
|
||||
* \brief Filters texture instructions.
|
||||
*
|
||||
* \param[in] instr NIR instruction.
|
||||
* \param[in] cb_data User callback data.
|
||||
* \return True if the instruction matches the filter.
|
||||
*/
|
||||
static bool is_tex(const nir_instr *instr, UNUSED const void *cb_data)
|
||||
{
|
||||
return instr->type == nir_instr_type_tex;
|
||||
nir_def_rewrite_uses(&tex->def, result);
|
||||
return true;
|
||||
}
|
||||
|
||||
/**
|
||||
|
|
@ -802,7 +798,10 @@ bool pco_nir_lower_tex(nir_shader *shader, pco_data *data, pco_ctx *ctx)
|
|||
.ctx = ctx,
|
||||
};
|
||||
|
||||
return nir_shader_lower_instructions(shader, is_tex, lower_tex, &state);
|
||||
return nir_shader_tex_pass(shader,
|
||||
lower_tex,
|
||||
nir_metadata_control_flow,
|
||||
&state);
|
||||
}
|
||||
|
||||
static enum util_format_type nir_type_to_util_type(nir_alu_type nir_type)
|
||||
|
|
@ -838,19 +837,15 @@ static enum pipe_format nir_type_to_pipe_format(nir_alu_type nir_type,
|
|||
pure_integer);
|
||||
}
|
||||
|
||||
static nir_def *lower_image(nir_builder *b, nir_instr *instr, void *cb_data)
|
||||
static bool
|
||||
lower_image(nir_builder *b, nir_intrinsic_instr *intr, void *cb_data)
|
||||
{
|
||||
nir_intrinsic_instr *intr = nir_instr_as_intrinsic(instr);
|
||||
struct state *state = cb_data;
|
||||
pco_data *data = state->data;
|
||||
pco_ctx *ctx = state->ctx;
|
||||
const struct pvr_device_info *dev_info = ctx->dev_info;
|
||||
enum glsl_sampler_dim image_dim = nir_intrinsic_image_dim(intr);
|
||||
bool is_array = nir_intrinsic_image_array(intr);
|
||||
enum pipe_format format = nir_intrinsic_format(intr);
|
||||
unsigned desc_set = nir_src_comp_as_uint(intr->src[0], 0);
|
||||
unsigned binding = nir_src_comp_as_uint(intr->src[0], 1);
|
||||
nir_def *elem = nir_channel(b, intr->src[0].ssa, 2);
|
||||
|
||||
b->cursor = nir_before_instr(&intr->instr);
|
||||
|
||||
nir_def *lod = NULL;
|
||||
switch (intr->intrinsic) {
|
||||
|
|
@ -872,9 +867,16 @@ static nir_def *lower_image(nir_builder *b, nir_instr *instr, void *cb_data)
|
|||
break;
|
||||
|
||||
default:
|
||||
UNREACHABLE("");
|
||||
return false;
|
||||
}
|
||||
|
||||
enum glsl_sampler_dim image_dim = nir_intrinsic_image_dim(intr);
|
||||
bool is_array = nir_intrinsic_image_array(intr);
|
||||
enum pipe_format format = nir_intrinsic_format(intr);
|
||||
unsigned desc_set = nir_src_comp_as_uint(intr->src[0], 0);
|
||||
unsigned binding = nir_src_comp_as_uint(intr->src[0], 1);
|
||||
nir_def *elem = nir_channel(b, intr->src[0].ssa, 2);
|
||||
|
||||
if (intr->intrinsic == nir_intrinsic_image_deref_size) {
|
||||
if (image_dim == GLSL_SAMPLER_DIM_BUF) {
|
||||
assert(intr->def.num_components == 1);
|
||||
|
|
@ -883,8 +885,11 @@ static nir_def *lower_image(nir_builder *b, nir_instr *instr, void *cb_data)
|
|||
elem,
|
||||
.desc_set = desc_set,
|
||||
.binding = binding);
|
||||
|
||||
return nir_channel(b, tex_meta, PCO_IMAGE_META_BUFFER_ELEMS);
|
||||
nir_def *buf_elems =
|
||||
nir_channel(b, tex_meta, PCO_IMAGE_META_BUFFER_ELEMS);
|
||||
nir_def_rewrite_uses(&intr->def, buf_elems);
|
||||
nir_instr_remove(&intr->instr);
|
||||
return true;
|
||||
}
|
||||
|
||||
nir_def *tex_state = nir_load_tex_state_pco(b,
|
||||
|
|
@ -907,7 +912,11 @@ static nir_def *lower_image(nir_builder *b, nir_instr *instr, void *cb_data)
|
|||
|
||||
data->common.uses.usclib = true;
|
||||
|
||||
return nir_trim_vector(b, size_comps, intr->def.num_components);
|
||||
nir_def *image_size =
|
||||
nir_trim_vector(b, size_comps, intr->def.num_components);
|
||||
nir_def_rewrite_uses(&intr->def, image_size);
|
||||
nir_instr_remove(&intr->instr);
|
||||
return true;
|
||||
}
|
||||
|
||||
nir_alu_type type = nir_type_invalid;
|
||||
|
|
@ -1071,22 +1080,24 @@ static nir_def *lower_image(nir_builder *b, nir_instr *instr, void *cb_data)
|
|||
|
||||
if (ia) {
|
||||
assert(!is_array);
|
||||
nir_load_const_instr *load =
|
||||
nir_def_as_load_const(intr->src[0].ssa);
|
||||
nir_load_const_instr *load = nir_def_as_load_const(intr->src[0].ssa);
|
||||
bool onchip = load->def.num_components == 4;
|
||||
|
||||
if (onchip) {
|
||||
unsigned ia_idx = nir_src_comp_as_uint(intr->src[0], 3);
|
||||
return nir_load_output(b,
|
||||
intr->def.num_components,
|
||||
intr->def.bit_size,
|
||||
nir_imm_int(b, 0),
|
||||
.base = ia_idx,
|
||||
.component = 0,
|
||||
.dest_type = nir_intrinsic_dest_type(intr),
|
||||
.io_semantics.location = FRAG_RESULT_COLOR,
|
||||
.io_semantics.num_slots = 1/*,
|
||||
.io_semantics.fb_fetch_output = true*/);
|
||||
nir_def *loaded_ia = nir_load_output(b,
|
||||
intr->def.num_components,
|
||||
intr->def.bit_size,
|
||||
nir_imm_int(b, 0),
|
||||
.base = ia_idx,
|
||||
.component = 0,
|
||||
.dest_type = nir_intrinsic_dest_type(intr),
|
||||
.io_semantics.location = FRAG_RESULT_COLOR,
|
||||
.io_semantics.num_slots = 1/*,
|
||||
.io_semantics.fb_fetch_output = true*/);
|
||||
nir_def_rewrite_uses(&intr->def, loaded_ia);
|
||||
nir_instr_remove(&intr->instr);
|
||||
return true;
|
||||
}
|
||||
}
|
||||
|
||||
|
|
@ -1210,10 +1221,13 @@ static nir_def *lower_image(nir_builder *b, nir_instr *instr, void *cb_data)
|
|||
|
||||
nir_def *addr_data = nir_vec4(b, addr_lo, addr_hi, compare, dma_data);
|
||||
|
||||
return nir_global_atomic_swap_pco(b,
|
||||
addr_data,
|
||||
.atomic_op =
|
||||
nir_intrinsic_atomic_op(intr));
|
||||
nir_def *atomic_swap = nir_global_atomic_swap_pco(
|
||||
b,
|
||||
addr_data,
|
||||
.atomic_op = nir_intrinsic_atomic_op(intr));
|
||||
nir_def_rewrite_uses(&intr->def, atomic_swap);
|
||||
nir_instr_remove(&intr->instr);
|
||||
return true;
|
||||
}
|
||||
|
||||
nir_def *dma_data = intr->src[3].ssa;
|
||||
|
|
@ -1222,9 +1236,13 @@ static nir_def *lower_image(nir_builder *b, nir_instr *instr, void *cb_data)
|
|||
|
||||
data->common.uses.usclib = true;
|
||||
|
||||
return nir_global_atomic_pco(b,
|
||||
addr_data,
|
||||
.atomic_op = nir_intrinsic_atomic_op(intr));
|
||||
nir_def *atomic =
|
||||
nir_global_atomic_pco(b,
|
||||
addr_data,
|
||||
.atomic_op = nir_intrinsic_atomic_op(intr));
|
||||
nir_def_rewrite_uses(&intr->def, atomic);
|
||||
nir_instr_remove(&intr->instr);
|
||||
return true;
|
||||
}
|
||||
|
||||
unsigned smp_desc = ia ? PCO_IA_SAMPLER : PCO_POINT_SAMPLER;
|
||||
|
|
@ -1383,31 +1401,14 @@ static nir_def *lower_image(nir_builder *b, nir_instr *instr, void *cb_data)
|
|||
|
||||
nir_intrinsic_instr *smp = pco_emit_nir_smp(b, ¶ms);
|
||||
|
||||
if (intr->intrinsic == nir_intrinsic_image_deref_load)
|
||||
return &smp->def;
|
||||
|
||||
return NIR_LOWER_INSTR_PROGRESS_REPLACE;
|
||||
}
|
||||
|
||||
static bool is_image(const nir_instr *instr, UNUSED const void *cb_data)
|
||||
{
|
||||
if (instr->type != nir_instr_type_intrinsic)
|
||||
return false;
|
||||
|
||||
nir_intrinsic_instr *intr = nir_instr_as_intrinsic(instr);
|
||||
switch (intr->intrinsic) {
|
||||
case nir_intrinsic_image_deref_load:
|
||||
case nir_intrinsic_image_deref_store:
|
||||
case nir_intrinsic_image_deref_atomic:
|
||||
case nir_intrinsic_image_deref_atomic_swap:
|
||||
case nir_intrinsic_image_deref_size:
|
||||
if (intr->intrinsic == nir_intrinsic_image_deref_load) {
|
||||
nir_def_rewrite_uses(&intr->def, &smp->def);
|
||||
nir_instr_remove(&intr->instr);
|
||||
return true;
|
||||
|
||||
default:
|
||||
break;
|
||||
}
|
||||
|
||||
return false;
|
||||
nir_instr_remove(&intr->instr);
|
||||
return true;
|
||||
}
|
||||
|
||||
bool pco_nir_lower_images(nir_shader *shader, pco_data *data, pco_ctx *ctx)
|
||||
|
|
@ -1416,5 +1417,9 @@ bool pco_nir_lower_images(nir_shader *shader, pco_data *data, pco_ctx *ctx)
|
|||
.data = data,
|
||||
.ctx = ctx,
|
||||
};
|
||||
return nir_shader_lower_instructions(shader, is_image, lower_image, &state);
|
||||
|
||||
return nir_shader_intrinsics_pass(shader,
|
||||
lower_image,
|
||||
nir_metadata_none,
|
||||
&state);
|
||||
}
|
||||
|
|
|
|||
|
|
@ -41,11 +41,11 @@ set_resource_used(pco_common_data *common, unsigned desc_set, unsigned binding)
|
|||
* \param[in] b NIR builder.
|
||||
* \param[in] intr NIR intrinsic instruction.
|
||||
* \param[in] common Shader common data.
|
||||
* \return The replacement/lowered def.
|
||||
* \return True if progress was made.
|
||||
*/
|
||||
static nir_def *lower_load_vulkan_descriptor(nir_builder *b,
|
||||
nir_intrinsic_instr *intr,
|
||||
pco_common_data *common)
|
||||
static bool lower_load_vulkan_descriptor(nir_builder *b,
|
||||
nir_intrinsic_instr *intr,
|
||||
pco_common_data *common)
|
||||
{
|
||||
nir_intrinsic_instr *vk_res_idx = nir_src_as_intrinsic(intr->src[0]);
|
||||
assert(vk_res_idx->intrinsic == nir_intrinsic_vulkan_resource_index);
|
||||
|
|
@ -58,8 +58,13 @@ static nir_def *lower_load_vulkan_descriptor(nir_builder *b,
|
|||
|
||||
set_resource_used(common, desc_set, binding);
|
||||
|
||||
b->cursor = nir_before_instr(&intr->instr);
|
||||
|
||||
uint32_t desc_set_binding = pco_pack_desc(desc_set, binding);
|
||||
return nir_imm_ivec3(b, desc_set_binding, elem, 0);
|
||||
nir_def *desc_ref = nir_imm_ivec3(b, desc_set_binding, elem, 0);
|
||||
nir_def_rewrite_uses(&intr->def, desc_ref);
|
||||
nir_instr_remove(&intr->instr);
|
||||
return true;
|
||||
}
|
||||
|
||||
static nir_def *array_elem_from_deref(nir_builder *b, nir_deref_instr *deref)
|
||||
|
|
@ -88,14 +93,13 @@ static inline bool is_comb_img_smp(unsigned desc_set,
|
|||
return binding_data->is_img_smp;
|
||||
}
|
||||
|
||||
static void lower_tex_deref_to_binding(nir_builder *b,
|
||||
static bool lower_tex_deref_to_binding(nir_builder *b,
|
||||
nir_tex_instr *tex,
|
||||
unsigned deref_index,
|
||||
pco_common_data *common)
|
||||
{
|
||||
nir_tex_src *deref_src = &tex->src[deref_index];
|
||||
nir_deref_instr *deref =
|
||||
nir_def_as_deref(deref_src->src.ssa);
|
||||
nir_deref_instr *deref = nir_def_as_deref(deref_src->src.ssa);
|
||||
|
||||
b->cursor = nir_before_instr(&tex->instr);
|
||||
|
||||
|
|
@ -117,9 +121,10 @@ static void lower_tex_deref_to_binding(nir_builder *b,
|
|||
}
|
||||
|
||||
nir_src_rewrite(&deref_src->src, elem);
|
||||
return true;
|
||||
}
|
||||
|
||||
static void
|
||||
static bool
|
||||
add_txf_sampler(nir_builder *b, nir_tex_instr *tex, pco_common_data *common)
|
||||
{
|
||||
int deref_index = nir_tex_instr_src_index(tex, nir_tex_src_backend1);
|
||||
|
|
@ -142,25 +147,29 @@ add_txf_sampler(nir_builder *b, nir_tex_instr *tex, pco_common_data *common)
|
|||
|
||||
tex->sampler_index = pco_pack_desc(desc_set, binding);
|
||||
nir_tex_instr_add_src(tex, nir_tex_src_backend2, elem);
|
||||
return true;
|
||||
}
|
||||
|
||||
static inline void
|
||||
static inline bool
|
||||
lower_tex_derefs(nir_builder *b, nir_tex_instr *tex, pco_common_data *common)
|
||||
{
|
||||
int deref_index;
|
||||
bool progress = false;
|
||||
|
||||
deref_index = nir_tex_instr_src_index(tex, nir_tex_src_texture_deref);
|
||||
if (deref_index >= 0)
|
||||
lower_tex_deref_to_binding(b, tex, deref_index, common);
|
||||
progress |= lower_tex_deref_to_binding(b, tex, deref_index, common);
|
||||
|
||||
deref_index = nir_tex_instr_src_index(tex, nir_tex_src_sampler_deref);
|
||||
if (deref_index >= 0)
|
||||
lower_tex_deref_to_binding(b, tex, deref_index, common);
|
||||
progress |= lower_tex_deref_to_binding(b, tex, deref_index, common);
|
||||
else if (tex->op == nir_texop_txf || tex->op == nir_texop_txf_ms)
|
||||
add_txf_sampler(b, tex, common);
|
||||
progress |= add_txf_sampler(b, tex, common);
|
||||
|
||||
return progress;
|
||||
}
|
||||
|
||||
static nir_def *
|
||||
static bool
|
||||
lower_image_derefs(nir_builder *b, nir_intrinsic_instr *intr, pco_data *data)
|
||||
{
|
||||
nir_src *deref_src = &intr->src[0];
|
||||
|
|
@ -200,7 +209,7 @@ lower_image_derefs(nir_builder *b, nir_intrinsic_instr *intr, pco_data *data)
|
|||
|
||||
nir_src_rewrite(deref_src, index);
|
||||
|
||||
return NIR_LOWER_INSTR_PROGRESS;
|
||||
return true;
|
||||
}
|
||||
|
||||
/* Sampler not needed for on-chip input attachments. */
|
||||
|
|
@ -217,18 +226,17 @@ lower_image_derefs(nir_builder *b, nir_intrinsic_instr *intr, pco_data *data)
|
|||
|
||||
nir_src_rewrite(deref_src, index);
|
||||
|
||||
return NIR_LOWER_INSTR_PROGRESS;
|
||||
return true;
|
||||
}
|
||||
|
||||
static nir_def *lower_is_null_descriptor(nir_builder *b,
|
||||
nir_intrinsic_instr *intr)
|
||||
static bool lower_is_null_descriptor(nir_builder *b, nir_intrinsic_instr *intr)
|
||||
{
|
||||
nir_src *deref_src = &intr->src[0];
|
||||
nir_deref_instr *deref = nir_src_as_deref(*deref_src);
|
||||
|
||||
/* Will be taken care of by lower_load_vulkan_descriptor. */
|
||||
if (!deref)
|
||||
return NULL;
|
||||
return false;
|
||||
|
||||
b->cursor = nir_before_instr(&intr->instr);
|
||||
|
||||
|
|
@ -242,102 +250,61 @@ static nir_def *lower_is_null_descriptor(nir_builder *b,
|
|||
nir_def *index = nir_vec2(b, nir_imm_int(b, desc_set_binding), elem);
|
||||
|
||||
nir_src_rewrite(deref_src, index);
|
||||
return NIR_LOWER_INSTR_PROGRESS;
|
||||
return true;
|
||||
}
|
||||
|
||||
/**
|
||||
* \brief Lowers a Vulkan-related instruction.
|
||||
* \brief Lowers a Vulkan-related texture instruction.
|
||||
*
|
||||
* \param[in] b NIR builder.
|
||||
* \param[in] instr NIR instruction.
|
||||
* \param[in] tex NIR texture instruction.
|
||||
* \param[in] cb_data User callback data.
|
||||
* \return The replacement/lowered def.
|
||||
* \return True if progress was made.
|
||||
*/
|
||||
static nir_def *lower_vk(nir_builder *b, nir_instr *instr, void *cb_data)
|
||||
static bool lower_vk_tex(nir_builder *b, nir_tex_instr *tex, void *cb_data)
|
||||
{
|
||||
pco_data *data = cb_data;
|
||||
pco_common_data *common = &data->common;
|
||||
|
||||
switch (instr->type) {
|
||||
case nir_instr_type_intrinsic: {
|
||||
nir_intrinsic_instr *intr = nir_instr_as_intrinsic(instr);
|
||||
switch (intr->intrinsic) {
|
||||
case nir_intrinsic_load_vulkan_descriptor:
|
||||
return lower_load_vulkan_descriptor(b, intr, common);
|
||||
|
||||
case nir_intrinsic_image_deref_load:
|
||||
case nir_intrinsic_image_deref_store:
|
||||
case nir_intrinsic_image_deref_atomic:
|
||||
case nir_intrinsic_image_deref_atomic_swap:
|
||||
case nir_intrinsic_image_deref_size:
|
||||
return lower_image_derefs(b, intr, data);
|
||||
|
||||
case nir_intrinsic_is_null_descriptor:
|
||||
return lower_is_null_descriptor(b, intr);
|
||||
|
||||
default:
|
||||
break;
|
||||
}
|
||||
|
||||
break;
|
||||
if (nir_tex_instr_src_index(tex, nir_tex_src_texture_deref) >= 0 ||
|
||||
nir_tex_instr_src_index(tex, nir_tex_src_sampler_deref) >= 0) {
|
||||
return lower_tex_derefs(b, tex, common);
|
||||
}
|
||||
|
||||
case nir_instr_type_tex: {
|
||||
nir_tex_instr *tex = nir_instr_as_tex(instr);
|
||||
lower_tex_derefs(b, tex, common);
|
||||
return NIR_LOWER_INSTR_PROGRESS;
|
||||
}
|
||||
|
||||
default:
|
||||
break;
|
||||
}
|
||||
|
||||
UNREACHABLE("");
|
||||
return false;
|
||||
}
|
||||
|
||||
/**
|
||||
* \brief Filters Vulkan-related instructions.
|
||||
* \brief Lowers a Vulkan-related intrinsic instruction.
|
||||
*
|
||||
* \param[in] instr NIR instruction.
|
||||
* \param[in] b NIR builder.
|
||||
* \param[in] intr NIR intrinsic instruction.
|
||||
* \param[in] cb_data User callback data.
|
||||
* \return True if the instruction matches the filter.
|
||||
* \return True if progress was made.
|
||||
*/
|
||||
static bool is_vk(const nir_instr *instr, UNUSED const void *cb_data)
|
||||
static bool
|
||||
lower_vk_intr(nir_builder *b, nir_intrinsic_instr *intr, void *cb_data)
|
||||
{
|
||||
switch (instr->type) {
|
||||
case nir_instr_type_intrinsic: {
|
||||
nir_intrinsic_instr *intr = nir_instr_as_intrinsic(instr);
|
||||
switch (intr->intrinsic) {
|
||||
case nir_intrinsic_load_vulkan_descriptor:
|
||||
case nir_intrinsic_is_null_descriptor:
|
||||
case nir_intrinsic_image_deref_load:
|
||||
case nir_intrinsic_image_deref_store:
|
||||
case nir_intrinsic_image_deref_atomic:
|
||||
case nir_intrinsic_image_deref_atomic_swap:
|
||||
case nir_intrinsic_image_deref_size:
|
||||
return true;
|
||||
pco_data *data = cb_data;
|
||||
pco_common_data *common = &data->common;
|
||||
|
||||
b->cursor = nir_before_instr(&intr->instr);
|
||||
|
||||
default:
|
||||
break;
|
||||
}
|
||||
switch (intr->intrinsic) {
|
||||
case nir_intrinsic_load_vulkan_descriptor:
|
||||
return lower_load_vulkan_descriptor(b, intr, common);
|
||||
|
||||
break;
|
||||
}
|
||||
case nir_intrinsic_image_deref_load:
|
||||
case nir_intrinsic_image_deref_store:
|
||||
case nir_intrinsic_image_deref_atomic:
|
||||
case nir_intrinsic_image_deref_atomic_swap:
|
||||
case nir_intrinsic_image_deref_size:
|
||||
return lower_image_derefs(b, intr, data);
|
||||
|
||||
case nir_instr_type_tex: {
|
||||
nir_tex_instr *tex = nir_instr_as_tex(instr);
|
||||
if (nir_tex_instr_src_index(tex, nir_tex_src_texture_deref) >= 0 ||
|
||||
nir_tex_instr_src_index(tex, nir_tex_src_sampler_deref) >= 0) {
|
||||
return true;
|
||||
}
|
||||
|
||||
FALLTHROUGH;
|
||||
}
|
||||
case nir_intrinsic_is_null_descriptor:
|
||||
return lower_is_null_descriptor(b, intr);
|
||||
|
||||
default:
|
||||
break;
|
||||
}
|
||||
|
||||
return false;
|
||||
}
|
||||
|
||||
|
|
@ -352,7 +319,14 @@ bool pco_nir_lower_vk(nir_shader *shader, pco_data *data)
|
|||
{
|
||||
bool progress = false;
|
||||
|
||||
progress |= nir_shader_lower_instructions(shader, is_vk, lower_vk, data);
|
||||
progress |= nir_shader_intrinsics_pass(shader,
|
||||
lower_vk_intr,
|
||||
nir_metadata_control_flow,
|
||||
data);
|
||||
progress |= nir_shader_tex_pass(shader,
|
||||
lower_vk_tex,
|
||||
nir_metadata_control_flow,
|
||||
data);
|
||||
|
||||
return progress;
|
||||
}
|
||||
|
|
|
|||
Loading…
Add table
Reference in a new issue