From 34bc9477de18a92e76ea7c536940a631323a83b6 Mon Sep 17 00:00:00 2001 From: Tony Wasserka Date: Wed, 7 Oct 2020 18:21:48 +0200 Subject: [PATCH] aco: Clean up symbol names and comments related to NGG MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Reviewed-by: Timur Kristóf Acked-by: Daniel Schürmann Part-of: --- src/amd/compiler/README.md | 8 ++--- src/amd/compiler/aco_assembler.cpp | 6 ++-- src/amd/compiler/aco_insert_exec_mask.cpp | 2 +- .../compiler/aco_instruction_selection.cpp | 36 +++++++++---------- .../aco_instruction_selection_setup.cpp | 22 ++++++------ src/amd/compiler/aco_ir.h | 18 +++++----- src/amd/compiler/aco_print_ir.cpp | 16 ++++----- src/amd/compiler/aco_scheduler.cpp | 2 +- 8 files changed, 55 insertions(+), 55 deletions(-) diff --git a/src/amd/compiler/README.md b/src/amd/compiler/README.md index 5a7cebbbb44..8fe2366ab2c 100644 --- a/src/amd/compiler/README.md +++ b/src/amd/compiler/README.md @@ -200,10 +200,10 @@ So, think about these as two independent shader programs slapped together. | GFX10/NGG HW stages: | LSHS | NGG GS | PS | ACO terminology | | -----------------------:|:----------|:-------------------|:---|:----------------| -| SW stages: only VS+PS: | | VS | FS | `ngg_vertex_gs`, `fragment_fs` | -| with tess: | VS + TCS | TES | FS | `vertex_tess_control_hs`, `ngg_tess_eval_gs`, `fragment_fs` | -| with GS: | | VS + GS | FS | `ngg_vertex_geometry_gs`, `fragment_fs` | -| with both: | VS + TCS | TES + GS | FS | `vertex_tess_control_hs`, `ngg_tess_eval_geometry_gs`, `fragment_fs` | +| SW stages: only VS+PS: | | VS | FS | `vertex_ngg`, `fragment_fs` | +| with tess: | VS + TCS | TES | FS | `vertex_tess_control_hs`, `tess_eval_ngg`, `fragment_fs` | +| with GS: | | VS + GS | FS | `vertex_geometry_ngg`, `fragment_fs` | +| with both: | VS + TCS | TES + GS | FS | `vertex_tess_control_hs`, `tess_eval_geometry_ngg`, `fragment_fs` | #### Compute pipeline diff --git a/src/amd/compiler/aco_assembler.cpp b/src/amd/compiler/aco_assembler.cpp index 24cb39e656c..7e678ad2c9a 100644 --- a/src/amd/compiler/aco_assembler.cpp +++ b/src/amd/compiler/aco_assembler.cpp @@ -732,7 +732,7 @@ void fix_exports(asm_context& ctx, std::vector& out, Program* program) { if ((*it)->format == Format::EXP) { Export_instruction* exp = static_cast((*it).get()); - if (program->stage.hw == HWStage::VS || program->stage.hw == HWStage::NGG_GS) { + if (program->stage.hw == HWStage::VS || program->stage.hw == HWStage::NGG) { if (exp->dest >= V_008DFC_SQ_EXP_POS && exp->dest <= (V_008DFC_SQ_EXP_POS + 3)) { exp->done = true; exported = true; @@ -752,7 +752,7 @@ void fix_exports(asm_context& ctx, std::vector& out, Program* program) if (!exported) { /* Abort in order to avoid a GPU hang. */ - bool is_vertex_or_ngg = (program->stage.hw == HWStage::VS || program->stage.hw == HWStage::NGG_GS); + bool is_vertex_or_ngg = (program->stage.hw == HWStage::VS || program->stage.hw == HWStage::NGG); aco_err(program, "Missing export in %s shader:", is_vertex_or_ngg ? "vertex or NGG" : "fragment"); aco_print_program(program, stderr); abort(); @@ -920,7 +920,7 @@ unsigned emit_program(Program* program, if (program->stage.hw == HWStage::VS || program->stage.hw == HWStage::FS || - program->stage.hw == HWStage::NGG_GS) + program->stage.hw == HWStage::NGG) fix_exports(ctx, code, program); for (Block& block : program->blocks) { diff --git a/src/amd/compiler/aco_insert_exec_mask.cpp b/src/amd/compiler/aco_insert_exec_mask.cpp index c6070e7531e..c5965c631a1 100644 --- a/src/amd/compiler/aco_insert_exec_mask.cpp +++ b/src/amd/compiler/aco_insert_exec_mask.cpp @@ -379,7 +379,7 @@ unsigned add_coupling_code(exec_ctx& ctx, Block* block, bld.insert(std::move(startpgm)); /* exec seems to need to be manually initialized with combined shaders */ - if (ctx.program->stage.num_sw_stages() > 1 || ctx.program->stage.hw == HWStage::NGG_GS) { + if (ctx.program->stage.num_sw_stages() > 1 || ctx.program->stage.hw == HWStage::NGG) { bld.sop1(Builder::s_mov, bld.exec(Definition(exec_mask)), bld.lm == s2 ? Operand(UINT64_MAX) : Operand(UINT32_MAX)); instructions[0]->definitions.pop_back(); } diff --git a/src/amd/compiler/aco_instruction_selection.cpp b/src/amd/compiler/aco_instruction_selection.cpp index babd769aa42..6c3969df433 100644 --- a/src/amd/compiler/aco_instruction_selection.cpp +++ b/src/amd/compiler/aco_instruction_selection.cpp @@ -4250,7 +4250,7 @@ void visit_store_ls_or_es_output(isel_context *ctx, nir_intrinsic_instr *instr) Temp lds_base; if (ctx->stage == vertex_geometry_gs || ctx->stage == tess_eval_geometry_gs || - ctx->stage == ngg_vertex_geometry_gs || ctx->stage == ngg_tess_eval_geometry_gs) { + ctx->stage == vertex_geometry_ngg || ctx->stage == tess_eval_geometry_ngg) { /* GFX9+: ES stage is merged into GS, data is passed between them using LDS. */ unsigned itemsize = ctx->stage.has(SWStage::VS) ? ctx->program->info->vs.es_info.esgs_itemsize @@ -4355,8 +4355,8 @@ void visit_store_output(isel_context *ctx, nir_intrinsic_instr *instr) if (ctx->stage == vertex_vs || ctx->stage == tess_eval_vs || ctx->stage == fragment_fs || - ctx->stage == ngg_vertex_gs || - ctx->stage == ngg_tess_eval_gs || + ctx->stage == vertex_ngg || + ctx->stage == tess_eval_ngg || ctx->shader->info.stage == MESA_SHADER_GEOMETRY) { bool stored_to_temps = store_output_to_temps(ctx, instr); if (!stored_to_temps) { @@ -8351,21 +8351,21 @@ void visit_intrinsic(isel_context *ctx, nir_intrinsic_instr *instr) break; } case nir_intrinsic_emit_vertex_with_counter: { - if (ctx->stage.hw == HWStage::NGG_GS) + if (ctx->stage.hw == HWStage::NGG) ngg_visit_emit_vertex_with_counter(ctx, instr); else visit_emit_vertex_with_counter(ctx, instr); break; } case nir_intrinsic_end_primitive_with_counter: { - if (ctx->stage.hw != HWStage::NGG_GS) { + if (ctx->stage.hw != HWStage::NGG) { unsigned stream = nir_intrinsic_stream_id(instr); bld.sopp(aco_opcode::s_sendmsg, bld.m0(ctx->gs_wave_id), -1, sendmsg_gs(true, false, stream)); } break; } case nir_intrinsic_set_vertex_and_primitive_count: { - if (ctx->stage.hw == HWStage::NGG_GS) + if (ctx->stage.hw == HWStage::NGG) ngg_visit_set_vertex_and_primitive_count(ctx, instr); /* unused in the legacy pipeline, the HW keeps track of this for us */ break; @@ -10082,7 +10082,7 @@ static bool visit_cf_list(isel_context *ctx, static void export_vs_varying(isel_context *ctx, int slot, bool is_pos, int *next_pos) { - assert(ctx->stage.hw == HWStage::VS || ctx->stage.hw == HWStage::NGG_GS); + assert(ctx->stage.hw == HWStage::VS || ctx->stage.hw == HWStage::NGG); int offset = (ctx->stage.has(SWStage::TES) && !ctx->stage.has(SWStage::GS)) ? ctx->program->info->tes.outinfo.vs_output_param_offset[slot] @@ -10179,13 +10179,13 @@ static void create_export_phis(isel_context *ctx) static void create_vs_exports(isel_context *ctx) { - assert(ctx->stage.hw == HWStage::VS || ctx->stage.hw == HWStage::NGG_GS); + assert(ctx->stage.hw == HWStage::VS || ctx->stage.hw == HWStage::NGG); radv_vs_output_info *outinfo = (ctx->stage.has(SWStage::TES) && !ctx->stage.has(SWStage::GS)) ? &ctx->program->info->tes.outinfo : &ctx->program->info->vs.outinfo; - if (outinfo->export_prim_id && ctx->stage.hw != HWStage::NGG_GS) { + if (outinfo->export_prim_id && ctx->stage.hw != HWStage::NGG) { ctx->outputs.mask[VARYING_SLOT_PRIMITIVE_ID] |= 0x1; if (ctx->stage.has(SWStage::TES)) ctx->outputs.temps[VARYING_SLOT_PRIMITIVE_ID * 4u] = get_arg(ctx, ctx->args->ac.tes_patch_id); @@ -11091,7 +11091,7 @@ Temp ngg_pack_prim_exp_arg(isel_context *ctx, unsigned num_vertices, const Temp Temp tmp; Temp gs_invocation_id; - if (ctx->stage == ngg_vertex_gs) + if (ctx->stage == vertex_ngg) gs_invocation_id = get_arg(ctx, ctx->args->ac.gs_invocation_id); for (unsigned i = 0; i < num_vertices; ++i) { @@ -11103,7 +11103,7 @@ Temp ngg_pack_prim_exp_arg(isel_context *ctx, unsigned num_vertices, const Temp tmp = vtxindex[i]; /* The initial edge flag is always false in tess eval shaders. */ - if (ctx->stage == ngg_vertex_gs) { + if (ctx->stage == vertex_ngg) { Temp edgeflag = bld.vop3(aco_opcode::v_bfe_u32, bld.def(v1), gs_invocation_id, Operand(8u + i), Operand(1u)); tmp = bld.vop3(aco_opcode::v_lshl_or_b32, bld.def(v1), edgeflag, Operand(10u * i + 9u), tmp); } @@ -11149,9 +11149,9 @@ void ngg_nogs_export_primitives(isel_context *ctx) assert(!ctx->stage.has(SWStage::GS)); - if (ctx->stage == ngg_vertex_gs) { + if (ctx->stage == vertex_ngg) { /* TODO: optimize for points & lines */ - } else if (ctx->stage == ngg_tess_eval_gs) { + } else if (ctx->stage == tess_eval_ngg) { if (ctx->shader->info.tess.point_mode) num_vertices_per_primitive = 1; else if (ctx->shader->info.tess.primitive_mode == GL_ISOLINES) @@ -11176,7 +11176,7 @@ void ngg_nogs_export_primitives(isel_context *ctx) ngg_emit_prim_export(ctx, num_vertices_per_primitive, vtxindex); /* Export primitive ID. */ - if (ctx->stage == ngg_vertex_gs && ctx->args->options->key.vs_common_out.export_prim_id) { + if (ctx->stage == vertex_ngg && ctx->args->options->key.vs_common_out.export_prim_id) { /* Copy Primitive IDs from GS threads to the LDS address corresponding to the ES thread of the provoking vertex. */ Temp prim_id = get_arg(ctx, ctx->args->ac.gs_prim_id); Temp provoking_vtx_index = vtxindex[0]; @@ -11201,7 +11201,7 @@ void ngg_nogs_export_vertices(isel_context *ctx) if (ctx->args->options->key.vs_common_out.export_prim_id) { Temp prim_id; - if (ctx->stage == ngg_vertex_gs) { + if (ctx->stage == vertex_ngg) { /* Wait for GS threads to store primitive ID in LDS. */ create_workgroup_barrier(bld); @@ -11211,7 +11211,7 @@ void ngg_nogs_export_vertices(isel_context *ctx) /* Load primitive ID from LDS. */ prim_id = load_lds(ctx, 4, bld.tmp(v1), addr, 0u, 4u); - } else if (ctx->stage == ngg_tess_eval_gs) { + } else if (ctx->stage == tess_eval_ngg) { /* TES: Just use the patch ID as the primitive ID. */ prim_id = get_arg(ctx, ctx->args->ac.tes_patch_id); } else { @@ -11650,8 +11650,8 @@ void select_program(Program *program, { isel_context ctx = setup_isel_context(program, shader_count, shaders, config, args, false); if_context ic_merged_wave_info; - bool ngg_no_gs = ctx.stage == ngg_vertex_gs || ctx.stage == ngg_tess_eval_gs; - bool ngg_gs = ctx.stage == ngg_vertex_geometry_gs || ctx.stage == ngg_tess_eval_geometry_gs; + bool ngg_no_gs = ctx.stage == vertex_ngg || ctx.stage == tess_eval_ngg; + bool ngg_gs = ctx.stage == vertex_geometry_ngg || ctx.stage == tess_eval_geometry_ngg; for (unsigned i = 0; i < shader_count; i++) { nir_shader *nir = shaders[i]; diff --git a/src/amd/compiler/aco_instruction_selection_setup.cpp b/src/amd/compiler/aco_instruction_selection_setup.cpp index c119c6eaeeb..c68b8f61fd2 100644 --- a/src/amd/compiler/aco_instruction_selection_setup.cpp +++ b/src/amd/compiler/aco_instruction_selection_setup.cpp @@ -429,13 +429,13 @@ setup_vs_output_info(isel_context *ctx, nir_shader *nir, void setup_vs_variables(isel_context *ctx, nir_shader *nir) { - if (ctx->stage == vertex_vs || ctx->stage == ngg_vertex_gs) { + if (ctx->stage == vertex_vs || ctx->stage == vertex_ngg) { radv_vs_output_info *outinfo = &ctx->program->info->vs.outinfo; setup_vs_output_info(ctx, nir, outinfo->export_prim_id, ctx->options->key.vs_common_out.export_clip_dists, outinfo); /* TODO: NGG streamout */ - if (ctx->stage.hw == HWStage::NGG_GS) + if (ctx->stage.hw == HWStage::NGG) assert(!ctx->args->shader_info->so.num_outputs); /* TODO: check if the shader writes edge flags (not in Vulkan) */ @@ -444,7 +444,7 @@ setup_vs_variables(isel_context *ctx, nir_shader *nir) ctx->tcs_num_inputs = ctx->program->info->vs.num_linked_outputs; } - if (ctx->stage == ngg_vertex_gs && ctx->args->options->key.vs_common_out.export_prim_id) { + if (ctx->stage == vertex_ngg && ctx->args->options->key.vs_common_out.export_prim_id) { /* We need to store the primitive IDs in LDS */ unsigned lds_size = ctx->program->info->ngg_info.esgs_ring_size; ctx->program->config->lds_size = (lds_size + ctx->program->lds_alloc_granule - 1) / @@ -456,7 +456,7 @@ void setup_gs_variables(isel_context *ctx, nir_shader *nir) { if (ctx->stage == vertex_geometry_gs || ctx->stage == tess_eval_geometry_gs) { ctx->program->config->lds_size = ctx->program->info->gs_ring_info.lds_size; /* Already in units of the alloc granularity */ - } else if (ctx->stage == ngg_vertex_geometry_gs || ctx->stage == ngg_tess_eval_geometry_gs) { + } else if (ctx->stage == vertex_geometry_ngg || ctx->stage == tess_eval_geometry_ngg) { radv_vs_output_info *outinfo = &ctx->program->info->vs.outinfo; setup_vs_output_info(ctx, nir, false, ctx->options->key.vs_common_out.export_clip_dists, outinfo); @@ -544,13 +544,13 @@ setup_tes_variables(isel_context *ctx, nir_shader *nir) ctx->tcs_num_patches = ctx->args->options->key.tes.num_patches; ctx->tcs_num_outputs = ctx->program->info->tes.num_linked_inputs; - if (ctx->stage == tess_eval_vs || ctx->stage == ngg_tess_eval_gs) { + if (ctx->stage == tess_eval_vs || ctx->stage == tess_eval_ngg) { radv_vs_output_info *outinfo = &ctx->program->info->tes.outinfo; setup_vs_output_info(ctx, nir, outinfo->export_prim_id, ctx->options->key.vs_common_out.export_clip_dists, outinfo); /* TODO: NGG streamout */ - if (ctx->stage.hw == HWStage::NGG_GS) + if (ctx->stage.hw == HWStage::NGG) assert(!ctx->args->shader_info->so.num_outputs); /* Tess eval shaders can't write edge flags, so this can be always true. */ @@ -1055,7 +1055,7 @@ setup_isel_context(Program* program, else if (sw_stage == SWStage::VS && !args->shader_info->vs.as_ls && !ngg) hw_stage = HWStage::VS; else if (sw_stage == SWStage::VS && ngg) - hw_stage = HWStage::NGG_GS; /* GFX10/NGG: VS without GS uses the HW GS stage */ + hw_stage = HWStage::NGG; /* GFX10/NGG: VS without GS uses the HW GS stage */ else if (sw_stage == SWStage::GS) hw_stage = HWStage::GS; else if (sw_stage == SWStage::FS) @@ -1067,7 +1067,7 @@ setup_isel_context(Program* program, else if (sw_stage == SWStage::VS_GS && gfx9_plus && !ngg) hw_stage = HWStage::GS; /* GFX6-9: VS+GS merged into a GS (and GFX10/legacy) */ else if (sw_stage == SWStage::VS_GS && ngg) - hw_stage = HWStage::NGG_GS; /* GFX10+: VS+GS merged into an NGG GS */ + hw_stage = HWStage::NGG; /* GFX10+: VS+GS merged into an NGG GS */ else if (sw_stage == SWStage::VS && args->shader_info->vs.as_ls) hw_stage = HWStage::LS; /* GFX6-8: VS is a Local Shader, when tessellation is used */ else if (sw_stage == SWStage::TCS) @@ -1077,13 +1077,13 @@ setup_isel_context(Program* program, else if (sw_stage == SWStage::TES && !args->shader_info->tes.as_es && !ngg) hw_stage = HWStage::VS; /* GFX6-9: TES without GS uses the HW VS stage (and GFX10/legacy) */ else if (sw_stage == SWStage::TES && !args->shader_info->tes.as_es && ngg) - hw_stage = HWStage::NGG_GS; /* GFX10/NGG: TES without GS uses the HW GS stage */ + hw_stage = HWStage::NGG; /* GFX10/NGG: TES without GS */ else if (sw_stage == SWStage::TES && args->shader_info->tes.as_es && !ngg) hw_stage = HWStage::ES; /* GFX6-8: TES is an Export Shader */ else if (sw_stage == SWStage::TES_GS && gfx9_plus && !ngg) hw_stage = HWStage::GS; /* GFX9: TES+GS merged into a GS (and GFX10/legacy) */ else if (sw_stage == SWStage::TES_GS && ngg) - hw_stage = HWStage::NGG_GS; /* GFX10+: TES+GS merged into an NGG GS */ + hw_stage = HWStage::NGG; /* GFX10+: TES+GS merged into an NGG GS */ else unreachable("Shader stage not implemented"); @@ -1126,7 +1126,7 @@ setup_isel_context(Program* program, /* Merged LSHS operates in workgroups, but can still have a different number of LS and HS invocations */ setup_tcs_info(&ctx, shaders[1], shaders[0]); program->workgroup_size = ctx.tcs_num_patches * MAX2(shaders[1]->info.tess.tcs_vertices_out, ctx.args->options->key.tcs.input_vertices); - } else if (program->stage.hw == HWStage::NGG_GS) { + } else if (program->stage.hw == HWStage::NGG) { gfx10_ngg_info &ngg_info = args->shader_info->ngg_info; /* Max ES (SW VS) threads */ diff --git a/src/amd/compiler/aco_ir.h b/src/amd/compiler/aco_ir.h index d17a4fb99b5..96c20c0b8aa 100644 --- a/src/amd/compiler/aco_ir.h +++ b/src/amd/compiler/aco_ir.h @@ -1532,11 +1532,11 @@ constexpr SWStage operator|(SWStage a, SWStage b) { */ enum class HWStage : uint8_t { VS, - ES, /* Export shader: pre-GS (VS or TES) on GFX6-8. Combined into GS on GFX9 (and GFX10/legacy). */ - GS, /* Geometry shader on GFX10/legacy and GFX6-9. */ - NGG_GS, /* Geometry shader on GFX10/NGG. */ - LS, /* Local shader: pre-TCS (VS) on GFX6-8. Combined into HS on GFX9 (and GFX10/legacy). */ - HS, /* Hull shader: TCS on GFX6-8. Merged VS and TCS on GFX9-10. */ + ES, /* Export shader: pre-GS (VS or TES) on GFX6-8. Combined into GS on GFX9 (and GFX10/legacy). */ + GS, /* Geometry shader on GFX10/legacy and GFX6-9. */ + NGG, /* Primitive shader, used to implement VS, TES, GS. */ + LS, /* Local shader: pre-TCS (VS) on GFX6-8. Combined into HS on GFX9 (and GFX10/legacy). */ + HS, /* Hull shader: TCS on GFX6-8. Merged VS and TCS on GFX9-10. */ FS, CS, }; @@ -1581,10 +1581,10 @@ static constexpr Stage compute_cs(HWStage::CS, SWStage::CS); static constexpr Stage tess_eval_vs(HWStage::VS, SWStage::TES); static constexpr Stage gs_copy_vs(HWStage::VS, SWStage::GSCopy); /* GFX10/NGG */ -static constexpr Stage ngg_vertex_gs(HWStage::NGG_GS, SWStage::VS); -static constexpr Stage ngg_vertex_geometry_gs(HWStage::NGG_GS, SWStage::VS_GS); -static constexpr Stage ngg_tess_eval_gs(HWStage::NGG_GS, SWStage::TES); -static constexpr Stage ngg_tess_eval_geometry_gs(HWStage::NGG_GS, SWStage::TES_GS); +static constexpr Stage vertex_ngg(HWStage::NGG, SWStage::VS); +static constexpr Stage vertex_geometry_ngg(HWStage::NGG, SWStage::VS_GS); +static constexpr Stage tess_eval_ngg(HWStage::NGG, SWStage::TES); +static constexpr Stage tess_eval_geometry_ngg(HWStage::NGG, SWStage::TES_GS); /* GFX9 (and GFX10 if NGG isn't used) */ static constexpr Stage vertex_geometry_gs(HWStage::GS, SWStage::VS_GS); static constexpr Stage vertex_tess_control_hs(HWStage::HS, SWStage::VS_TCS); diff --git a/src/amd/compiler/aco_print_ir.cpp b/src/amd/compiler/aco_print_ir.cpp index 8b8b5d0f306..57435c2e433 100644 --- a/src/amd/compiler/aco_print_ir.cpp +++ b/src/amd/compiler/aco_print_ir.cpp @@ -830,14 +830,14 @@ static void print_stage(Stage stage, FILE *output) fprintf(output, "vertex_geometry_gs"); else if (stage == tess_eval_geometry_gs) fprintf(output, "tess_eval_geometry_gs"); - else if (stage == ngg_vertex_gs) - fprintf(output, "ngg_vertex_gs"); - else if (stage == ngg_tess_eval_gs) - fprintf(output, "ngg_tess_eval_gs"); - else if (stage == ngg_vertex_geometry_gs) - fprintf(output, "ngg_vertex_geometry_gs"); - else if (stage == ngg_tess_eval_geometry_gs) - fprintf(output, "ngg_tess_eval_geometry_gs"); + else if (stage == vertex_ngg) + fprintf(output, "vertex_ngg"); + else if (stage == tess_eval_ngg) + fprintf(output, "tess_eval_ngg"); + else if (stage == vertex_geometry_ngg) + fprintf(output, "vertex_geometry_ngg"); + else if (stage == tess_eval_geometry_ngg) + fprintf(output, "tess_eval_geometry_ngg"); else fprintf(output, "unknown"); diff --git a/src/amd/compiler/aco_scheduler.cpp b/src/amd/compiler/aco_scheduler.cpp index 83e5dcf0b80..1ad01cb37cc 100644 --- a/src/amd/compiler/aco_scheduler.cpp +++ b/src/amd/compiler/aco_scheduler.cpp @@ -862,7 +862,7 @@ void schedule_block(sched_ctx& ctx, Program *program, Block* block, live& live_v } } - if ((program->stage.hw == HWStage::VS || program->stage.hw == HWStage::NGG_GS) && + if ((program->stage.hw == HWStage::VS || program->stage.hw == HWStage::NGG) && (block->kind & block_kind_export_end)) { /* Try to move position exports as far up as possible, to reduce register * usage and because ISA reference guides say so. */