Module: Mesa Branch: main Commit: 028d0590f85c418cf470510cafcbe8026c2c4208 URL: http://cgit.freedesktop.org/mesa/mesa/commit/?id=028d0590f85c418cf470510cafcbe8026c2c4208
Author: Qiang Yu <[email protected]> Date: Sun Jun 12 20:36:39 2022 +0800 radeonsi: replace llvm ngg vs/tes with nir lowering Acked-by: Pierre-Eric Pelloux-Prayer <[email protected]> Reviewed-by: Marek Olšák <[email protected]> Signed-off-by: Qiang Yu <[email protected]> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/17109> --- src/gallium/drivers/radeonsi/gfx10_shader_ngg.c | 66 +++++++- src/gallium/drivers/radeonsi/si_shader.c | 122 +++++++++------ src/gallium/drivers/radeonsi/si_shader.h | 3 +- src/gallium/drivers/radeonsi/si_shader_internal.h | 14 +- src/gallium/drivers/radeonsi/si_shader_llvm.c | 175 +++++++--------------- src/gallium/drivers/radeonsi/si_shader_llvm_gs.c | 2 +- src/gallium/drivers/radeonsi/si_shader_llvm_vs.c | 18 +-- src/gallium/drivers/radeonsi/si_shader_nir.c | 21 +-- 8 files changed, 215 insertions(+), 206 deletions(-) diff --git a/src/gallium/drivers/radeonsi/gfx10_shader_ngg.c b/src/gallium/drivers/radeonsi/gfx10_shader_ngg.c index 8ecfd74e24f..c2dc9d4fd6f 100644 --- a/src/gallium/drivers/radeonsi/gfx10_shader_ngg.c +++ b/src/gallium/drivers/radeonsi/gfx10_shader_ngg.c @@ -119,6 +119,37 @@ static LLVMValueRef ngg_get_vertices_per_prim(struct si_shader_context *ctx, uns } } +unsigned gfx10_ngg_get_vertices_per_prim(struct si_shader *shader) +{ + const struct si_shader_info *info = &shader->selector->info; + + if (shader->selector->stage == MESA_SHADER_GEOMETRY) + return u_vertices_per_prim(info->base.gs.output_primitive); + else if (shader->selector->stage == MESA_SHADER_VERTEX) { + if (info->base.vs.blit_sgprs_amd) { + /* Blits always use axis-aligned rectangles with 3 vertices. */ + return 3; + } else if (shader->key.ge.opt.ngg_culling & SI_NGG_CULL_LINES) + return 2; + else { + /* We always build up all three indices for the prim export + * independent of the primitive type. The additional garbage + * data shouldn't hurt. This is used by exports and streamout. + */ + return 3; + } + } else { + assert(shader->selector->stage == MESA_SHADER_TESS_EVAL); + + if (info->base.tess.point_mode) + return 1; + else if (info->base.tess._primitive_mode == TESS_PRIMITIVE_ISOLINES) + return 2; + else + return 3; + } +} + bool gfx10_ngg_export_prim_early(struct si_shader *shader) { struct si_shader_selector *sel = shader->selector; @@ -2398,11 +2429,17 @@ static void clamp_gsprims_to_esverts(unsigned *max_gsprims, unsigned max_esverts unsigned gfx10_ngg_get_scratch_dw_size(struct si_shader *shader) { const struct si_shader_selector *sel = shader->selector; + bool uses_streamout = si_shader_uses_streamout(shader); - if (sel->stage == MESA_SHADER_GEOMETRY && si_shader_uses_streamout(shader)) - return 44; - - return 8; + if (sel->stage == MESA_SHADER_GEOMETRY) { + return uses_streamout ? 44 : 8; + } else { + return ac_ngg_get_scratch_lds_size(sel->stage, + si_get_max_workgroup_size(shader), + shader->wave_size, + uses_streamout, + shader->key.ge.opt.ngg_culling) / 4; + } } /** @@ -2469,8 +2506,25 @@ retry_select_mode: } } else { /* VS and TES. */ - /* LDS size for passing data from ES to GS. */ - esvert_lds_size = ngg_nogs_vertex_size(shader); + + bool uses_instance_id = gs_sel->info.uses_instanceid; + bool uses_primitive_id = gs_sel->info.uses_primid; + if (gs_stage == MESA_SHADER_VERTEX) { + uses_instance_id |= + shader->key.ge.part.vs.prolog.instance_divisor_is_one || + shader->key.ge.part.vs.prolog.instance_divisor_is_fetched; + } else { + uses_primitive_id |= shader->key.ge.mono.u.vs_export_prim_id; + } + + esvert_lds_size = ac_ngg_nogs_get_pervertex_lds_size( + gs_stage, gs_sel->info.num_outputs, + si_shader_uses_streamout(shader), + shader->key.ge.mono.u.vs_export_prim_id, + gfx10_ngg_writes_user_edgeflags(shader), + shader->key.ge.opt.ngg_culling, + uses_instance_id, + uses_primitive_id) / 4; } unsigned max_gsprims = max_gsprims_base; diff --git a/src/gallium/drivers/radeonsi/si_shader.c b/src/gallium/drivers/radeonsi/si_shader.c index 73dd8d4ee0d..ef864365511 100644 --- a/src/gallium/drivers/radeonsi/si_shader.c +++ b/src/gallium/drivers/radeonsi/si_shader.c @@ -227,7 +227,7 @@ unsigned si_get_max_workgroup_size(const struct si_shader *shader) switch (shader->selector->stage) { case MESA_SHADER_VERTEX: case MESA_SHADER_TESS_EVAL: - return shader->key.ge.as_ngg ? 128 : 0; + return shader->key.ge.as_ngg ? shader->selector->screen->ngg_subgroup_size : 0; case MESA_SHADER_TESS_CTRL: /* Return this so that LLVM doesn't remove s_barrier @@ -397,7 +397,7 @@ void si_add_arg_checked(struct ac_shader_args *args, enum ac_arg_regfile file, u ac_add_arg(args, file, registers, type, arg); } -void si_init_shader_args(struct si_shader_context *ctx, bool ngg_cull_shader) +void si_init_shader_args(struct si_shader_context *ctx) { struct si_shader *shader = ctx->shader; unsigned i, num_returns, num_return_sgprs; @@ -613,36 +613,12 @@ void si_init_shader_args(struct si_shader_context *ctx, bool ngg_cull_shader) declare_tes_input_vgprs(ctx); } - if ((ctx->shader->key.ge.as_es || ngg_cull_shader) && + if (ctx->shader->key.ge.as_es && (ctx->stage == MESA_SHADER_VERTEX || ctx->stage == MESA_SHADER_TESS_EVAL)) { - unsigned num_user_sgprs, num_vgprs; - - if (ctx->stage == MESA_SHADER_VERTEX && ngg_cull_shader) { - /* For the NGG cull shader, add 1 SGPR to hold - * the vertex buffer pointer. - */ - num_user_sgprs = GFX9_GS_NUM_USER_SGPR + 1; - - if (shader->selector->info.num_vbos_in_user_sgprs) { - assert(num_user_sgprs <= SI_SGPR_VS_VB_DESCRIPTOR_FIRST); - num_user_sgprs = - SI_SGPR_VS_VB_DESCRIPTOR_FIRST + shader->selector->info.num_vbos_in_user_sgprs * 4; - } - } else { - num_user_sgprs = GFX9_GS_NUM_USER_SGPR; - } - - /* The NGG cull shader has to return all 9 VGPRs. - * - * The normal merged ESGS shader only has to return the 5 VGPRs - * for the GS stage. - */ - num_vgprs = ngg_cull_shader ? 9 : 5; - /* ES return values are inputs to GS. */ - for (i = 0; i < 8 + num_user_sgprs; i++) + for (i = 0; i < 8 + GFX9_GS_NUM_USER_SGPR; i++) ac_add_return(&ctx->args, AC_ARG_SGPR); - for (i = 0; i < num_vgprs; i++) + for (i = 0; i < 5; i++) ac_add_return(&ctx->args, AC_ARG_VGPR); } break; @@ -1403,17 +1379,13 @@ static void si_dump_shader_key(const struct si_shader *shader, FILE *f) } bool si_vs_needs_prolog(const struct si_shader_selector *sel, - const struct si_vs_prolog_bits *prolog_key, - const union si_shader_key *key, bool ngg_cull_shader, - bool is_gs) + const struct si_vs_prolog_bits *prolog_key) { assert(sel->stage == MESA_SHADER_VERTEX); /* VGPR initialization fixup for Vega10 and Raven is always done in the * VS prolog. */ - return sel->info.vs_needs_prolog || prolog_key->ls_vgpr_fix || - /* The 2nd VS prolog loads input VGPRs from LDS */ - (key->ge.opt.ngg_culling && !ngg_cull_shader && !is_gs); + return sel->info.vs_needs_prolog || prolog_key->ls_vgpr_fix; } /** @@ -1422,13 +1394,12 @@ bool si_vs_needs_prolog(const struct si_shader_selector *sel, * * \param info Shader info of the vertex shader. * \param num_input_sgprs Number of input SGPRs for the vertex shader. - * \param has_old_ Whether the preceding shader part is the NGG cull shader. * \param prolog_key Key of the VS prolog * \param shader_out The vertex shader, or the next shader if merging LS+HS or ES+GS. * \param key Output shader part key. */ void si_get_vs_prolog_key(const struct si_shader_info *info, unsigned num_input_sgprs, - bool ngg_cull_shader, const struct si_vs_prolog_bits *prolog_key, + const struct si_vs_prolog_bits *prolog_key, struct si_shader *shader_out, union si_shader_part_key *key) { memset(key, 0, sizeof(*key)); @@ -1440,10 +1411,6 @@ void si_get_vs_prolog_key(const struct si_shader_info *info, unsigned num_input_ key->vs_prolog.as_es = shader_out->key.ge.as_es; key->vs_prolog.as_ngg = shader_out->key.ge.as_ngg; - if (shader_out->selector->stage != MESA_SHADER_GEOMETRY && - !ngg_cull_shader && shader_out->key.ge.opt.ngg_culling) - key->vs_prolog.load_vgprs_after_culling = 1; - if (shader_out->selector->stage == MESA_SHADER_TESS_CTRL) { key->vs_prolog.as_ls = 1; key->vs_prolog.num_merged_next_stage_vgprs = 2; @@ -1647,6 +1614,68 @@ static bool si_lower_io_to_mem(struct si_shader *shader, nir_shader *nir, return false; } +static void si_lower_ngg(struct si_shader *shader, nir_shader *nir) +{ + struct si_shader_selector *sel = shader->selector; + const union si_shader_key *key = &shader->key; + assert(key->ge.as_ngg); + + ac_nir_lower_ngg_options options = { + .family = sel->screen->info.family, + .gfx_level = sel->screen->info.gfx_level, + .max_workgroup_size = si_get_max_workgroup_size(shader), + .wave_size = shader->wave_size, + .can_cull = !!key->ge.opt.ngg_culling, + .disable_streamout = key->ge.opt.remove_streamout, + .vs_output_param_offset = shader->info.vs_output_param_offset, + }; + + if (nir->info.stage == MESA_SHADER_VERTEX || + nir->info.stage == MESA_SHADER_TESS_EVAL) { + /* Per instance inputs, used to remove instance load after culling. */ + unsigned instance_rate_inputs = 0; + + if (nir->info.stage == MESA_SHADER_VERTEX) { + instance_rate_inputs = + key->ge.part.vs.prolog.instance_divisor_is_one | + key->ge.part.vs.prolog.instance_divisor_is_fetched; + + /* Manually mark the instance ID used, so the shader can repack it. */ + if (instance_rate_inputs) + BITSET_SET(nir->info.system_values_read, SYSTEM_VALUE_INSTANCE_ID); + } else { + /* Manually mark the primitive ID used, so the shader can repack it. */ + if (key->ge.mono.u.vs_export_prim_id) + BITSET_SET(nir->info.system_values_read, SYSTEM_VALUE_PRIMITIVE_ID); + } + + unsigned clip_plane_enable = + SI_NGG_CULL_GET_CLIP_PLANE_ENABLE(key->ge.opt.ngg_culling); + unsigned clipdist_mask = + (sel->info.clipdist_mask & clip_plane_enable) | sel->info.culldist_mask; + + options.num_vertices_per_primitive = gfx10_ngg_get_vertices_per_prim(shader); + options.early_prim_export = gfx10_ngg_export_prim_early(shader); + options.passthrough = gfx10_is_ngg_passthrough(shader); + options.use_edgeflags = gfx10_edgeflags_have_effect(shader); + options.has_gen_prim_query = options.has_xfb_prim_query = + sel->screen->use_ngg_streamout && !sel->info.base.vs.blit_sgprs_amd; + options.primitive_id_location = + key->ge.mono.u.vs_export_prim_id ? sel->info.num_outputs : -1; + options.instance_rate_inputs = instance_rate_inputs; + options.clipdist_enable_mask = clipdist_mask; + options.user_clip_plane_enable_mask = clip_plane_enable; + + NIR_PASS_V(nir, ac_nir_lower_ngg_nogs, &options); + } + + /* may generate some subgroup op like ballot */ + NIR_PASS_V(nir, nir_lower_subgroups, &si_nir_subgroups_options); + + /* may generate some vector output store */ + NIR_PASS_V(nir, nir_lower_io_to_scalar, nir_var_shader_out); +} + struct nir_shader *si_deserialize_shader(struct si_shader_selector *sel) { struct pipe_screen *screen = &sel->screen->b; @@ -1878,6 +1907,12 @@ struct nir_shader *si_get_nir_shader(struct si_shader *shader, bool *free_nir, if (is_last_vgt_stage) si_assign_param_offsets(nir, shader); + /* Only lower last VGT NGG shader stage. */ + if (sel->stage < MESA_SHADER_GEOMETRY && key->ge.as_ngg && !key->ge.as_es) { + si_lower_ngg(shader, nir); + opt_offsets = true; + } + if (progress2 || opt_offsets) si_nir_opts(sel->screen, nir, false); @@ -2176,13 +2211,12 @@ static bool si_get_vs_prolog(struct si_screen *sscreen, struct ac_llvm_compiler { struct si_shader_selector *vs = main_part->selector; - if (!si_vs_needs_prolog(vs, key, &shader->key, false, - shader->selector->stage == MESA_SHADER_GEOMETRY)) + if (!si_vs_needs_prolog(vs, key)) return true; /* Get the prolog. */ union si_shader_part_key prolog_key; - si_get_vs_prolog_key(&vs->info, main_part->info.num_input_sgprs, false, key, shader, + si_get_vs_prolog_key(&vs->info, main_part->info.num_input_sgprs, key, shader, &prolog_key); shader->prolog = diff --git a/src/gallium/drivers/radeonsi/si_shader.h b/src/gallium/drivers/radeonsi/si_shader.h index b1ff7fe654d..83cad1a8e21 100644 --- a/src/gallium/drivers/radeonsi/si_shader.h +++ b/src/gallium/drivers/radeonsi/si_shader.h @@ -613,7 +613,6 @@ union si_shader_part_key { unsigned as_ls : 1; unsigned as_es : 1; unsigned as_ngg : 1; - unsigned load_vgprs_after_culling : 1; /* Prologs for monolithic shaders shouldn't set EXEC. */ unsigned is_monolithic : 1; } vs_prolog; @@ -1002,6 +1001,8 @@ struct si_shader *si_generate_gs_copy_shader(struct si_screen *sscreen, struct util_debug_callback *debug); /* si_shader_nir.c */ +extern const nir_lower_subgroups_options si_nir_subgroups_options; + void si_nir_opts(struct si_screen *sscreen, struct nir_shader *nir, bool first); void si_nir_late_opts(nir_shader *nir); char *si_finalize_nir(struct pipe_screen *screen, void *nirptr); diff --git a/src/gallium/drivers/radeonsi/si_shader_internal.h b/src/gallium/drivers/radeonsi/si_shader_internal.h index 405deed969b..164c20a55be 100644 --- a/src/gallium/drivers/radeonsi/si_shader_internal.h +++ b/src/gallium/drivers/radeonsi/si_shader_internal.h @@ -161,13 +161,12 @@ bool si_is_multi_part_shader(struct si_shader *shader); bool si_is_merged_shader(struct si_shader *shader); void si_add_arg_checked(struct ac_shader_args *args, enum ac_arg_regfile file, unsigned registers, enum ac_arg_type type, struct ac_arg *arg, unsigned idx); -void si_init_shader_args(struct si_shader_context *ctx, bool ngg_cull_shader); +void si_init_shader_args(struct si_shader_context *ctx); unsigned si_get_max_workgroup_size(const struct si_shader *shader); bool si_vs_needs_prolog(const struct si_shader_selector *sel, - const struct si_vs_prolog_bits *prolog_key, - const union si_shader_key *key, bool ngg_cull_shader, bool is_gs); + const struct si_vs_prolog_bits *prolog_key); void si_get_vs_prolog_key(const struct si_shader_info *info, unsigned num_input_sgprs, - bool ngg_cull_shader, const struct si_vs_prolog_bits *prolog_key, + const struct si_vs_prolog_bits *prolog_key, struct si_shader *shader_out, union si_shader_part_key *key); struct nir_shader *si_get_nir_shader(struct si_shader *shader, bool *free_nir, uint64_t tcs_vgpr_only_inputs); @@ -180,6 +179,7 @@ void si_fix_resource_usage(struct si_screen *sscreen, struct si_shader *shader); /* gfx10_shader_ngg.c */ LLVMValueRef gfx10_get_thread_id_in_tg(struct si_shader_context *ctx); +unsigned gfx10_ngg_get_vertices_per_prim(struct si_shader *shader); bool gfx10_ngg_export_prim_early(struct si_shader *shader); void gfx10_ngg_build_sendmsg_gs_alloc_req(struct si_shader_context *ctx); void gfx10_ngg_build_export_prim(struct si_shader_context *ctx, LLVMValueRef user_edgeflags[3], @@ -205,7 +205,7 @@ void si_llvm_context_init(struct si_shader_context *ctx, struct si_screen *sscre struct ac_llvm_compiler *compiler, unsigned wave_size); void si_llvm_create_func(struct si_shader_context *ctx, const char *name, LLVMTypeRef *return_types, unsigned num_return_elems, unsigned max_workgroup_size); -void si_llvm_create_main_func(struct si_shader_context *ctx, bool ngg_cull_shader); +void si_llvm_create_main_func(struct si_shader_context *ctx); void si_llvm_optimize_module(struct si_shader_context *ctx); void si_llvm_dispose(struct si_shader_context *ctx); LLVMValueRef si_buffer_load_const(struct si_shader_context *ctx, LLVMValueRef resource, @@ -228,7 +228,7 @@ void si_build_wrapper_function(struct si_shader_context *ctx, struct ac_llvm_poi enum ac_arg_type *main_arg_types, bool same_thread_count); bool si_llvm_translate_nir(struct si_shader_context *ctx, struct si_shader *shader, - struct nir_shader *nir, bool free_nir, bool ngg_cull_shader); + struct nir_shader *nir, bool free_nir); bool si_llvm_compile_shader(struct si_screen *sscreen, struct ac_llvm_compiler *compiler, struct si_shader *shader, const struct pipe_stream_output_info *so, struct util_debug_callback *debug, struct nir_shader *nir, @@ -278,6 +278,6 @@ void si_llvm_build_vs_exports(struct si_shader_context *ctx, LLVMValueRef num_ex struct si_shader_output_values *outputs, unsigned noutput); void si_llvm_vs_build_end(struct si_shader_context *ctx); void si_llvm_build_vs_prolog(struct si_shader_context *ctx, union si_shader_part_key *key); -void si_llvm_init_vs_callbacks(struct si_shader_context *ctx, bool ngg_cull_shader); +void si_llvm_init_vs_callbacks(struct si_shader_context *ctx); #endif diff --git a/src/gallium/drivers/radeonsi/si_shader_llvm.c b/src/gallium/drivers/radeonsi/si_shader_llvm.c index 94abe102a11..07d47691e16 100644 --- a/src/gallium/drivers/radeonsi/si_shader_llvm.c +++ b/src/gallium/drivers/radeonsi/si_shader_llvm.c @@ -197,21 +197,21 @@ void si_llvm_create_func(struct si_shader_context *ctx, const char *name, LLVMTy ac_llvm_set_target_features(ctx->main_fn.value, &ctx->ac); } -void si_llvm_create_main_func(struct si_shader_context *ctx, bool ngg_cull_shader) +void si_llvm_create_main_func(struct si_shader_context *ctx) { struct si_shader *shader = ctx->shader; LLVMTypeRef returns[AC_MAX_ARGS]; unsigned i; - si_init_shader_args(ctx, ngg_cull_shader); + si_init_shader_args(ctx); for (i = 0; i < ctx->args.num_sgprs_returned; i++) returns[i] = ctx->ac.i32; /* SGPR */ for (; i < ctx->args.return_count; i++) returns[i] = ctx->ac.f32; /* VGPR */ - si_llvm_create_func(ctx, ngg_cull_shader ? "ngg_cull_main" : "main", returns, - ctx->args.return_count, si_get_max_workgroup_size(shader)); + si_llvm_create_func(ctx, "main", returns, ctx->args.return_count, + si_get_max_workgroup_size(shader)); /* Reserve register locations for VGPR inputs the PS prolog may need. */ if (ctx->stage == MESA_SHADER_FRAGMENT && !ctx->shader->is_monolithic) { @@ -954,7 +954,7 @@ static LLVMValueRef si_llvm_load_streamout_buffer(struct ac_shader_abi *abi, uns } bool si_llvm_translate_nir(struct si_shader_context *ctx, struct si_shader *shader, - struct nir_shader *nir, bool free_nir, bool ngg_cull_shader) + struct nir_shader *nir, bool free_nir) { struct si_shader_selector *sel = shader->selector; const struct si_shader_info *info = &sel->info; @@ -975,7 +975,7 @@ bool si_llvm_translate_nir(struct si_shader_context *ctx, struct si_shader *shad ctx->abi.atomic_add_prim_count = gfx10_ngg_atomic_add_prim_count; si_llvm_init_resource_callbacks(ctx); - si_llvm_create_main_func(ctx, ngg_cull_shader); + si_llvm_create_main_func(ctx); if (ctx->stage <= MESA_SHADER_GEOMETRY && (ctx->shader->key.ge.as_es || ctx->stage == MESA_SHADER_GEOMETRY)) @@ -983,7 +983,7 @@ bool si_llvm_translate_nir(struct si_shader_context *ctx, struct si_shader *shad switch (ctx->stage) { case MESA_SHADER_VERTEX: - si_llvm_init_vs_callbacks(ctx, ngg_cull_shader); + si_llvm_init_vs_callbacks(ctx); /* preload instance_divisor_constbuf to be used for input load after culling */ if (ctx->shader->key.ge.opt.ngg_culling && @@ -1100,56 +1100,49 @@ bool si_llvm_translate_nir(struct si_shader_context *ctx, struct si_shader *shad * determined during linking / PM4 creation. */ si_llvm_declare_esgs_ring(ctx); + ctx->ac.lds.value = ctx->esgs_ring; + ctx->ac.lds.pointee_type = ctx->ac.i32; /* This is really only needed when streamout and / or vertex * compaction is enabled. */ - if (!ctx->gs_ngg_scratch.value && (ctx->so.num_outputs || shader->key.ge.opt.ngg_culling)) { + if (si_shader_uses_streamout(shader) || shader->key.ge.opt.ngg_culling) { LLVMTypeRef asi32 = LLVMArrayType(ctx->ac.i32, gfx10_ngg_get_scratch_dw_size(shader)); ctx->gs_ngg_scratch = (struct ac_llvm_pointer) { .value = LLVMAddGlobalInAddressSpace(ctx->ac.module, asi32, "ngg_scratch", AC_ADDR_SPACE_LDS), .pointee_type = asi32 }; LLVMSetInitializer(ctx->gs_ngg_scratch.value, LLVMGetUndef(asi32)); - LLVMSetAlignment(ctx->gs_ngg_scratch.value, 4); + LLVMSetAlignment(ctx->gs_ngg_scratch.value, 8); } } /* For merged shaders (VS-TCS, VS-GS, TES-GS): */ if (ctx->screen->info.gfx_level >= GFX9 && si_is_merged_shader(shader)) { - /* TES is special because it has only 1 shader part if NGG shader culling is disabled, - * and therefore it doesn't use the wrapper function. + /* Set EXEC = ~0 before the first shader. For monolithic shaders, the wrapper + * function does this. */ - bool no_wrapper_func = ctx->stage == MESA_SHADER_TESS_EVAL && !shader->key.ge.as_es && - !shader->key.ge.opt.ngg_culling; + if (ctx->stage == MESA_SHADER_TESS_EVAL) { + /* TES has only 1 shader part, therefore it doesn't use the wrapper function. */ + if (!shader->is_monolithic || !shader->key.ge.as_es) + ac_init_exec_full_mask(&ctx->ac); + } else if (ctx->stage == MESA_SHADER_VERTEX) { + /* If the prolog is present, EXEC is set there instead. */ + if (!si_vs_needs_prolog(sel, &shader->key.ge.part.vs.prolog)) { + /* When no prolog, only mono VS with TCS/GS present has wrapper function. */ + if (!(shader->is_monolithic && (shader->key.ge.as_ls || shader->key.ge.as_es))) + ac_init_exec_full_mask(&ctx->ac); + } + } - /* Set EXEC = ~0 before the first shader. If the prolog is present, EXEC is set there - * instead. For monolithic shaders, the wrapper function does this. - */ - if ((!shader->is_monolithic || no_wrapper_func) && - (ctx->stage == MESA_SHADER_TESS_EVAL || - (ctx->stage == MESA_SHADER_VERTEX && - !si_vs_needs_prolog(sel, &shader->key.ge.part.vs.prolog, &shader->key, ngg_cull_shader, - false)))) - ac_init_exec_full_mask(&ctx->ac); - - /* NGG VS and NGG TES: Send gs_alloc_req and the prim export at the beginning to decrease - * register usage. + /* NGG VS and NGG TES: nir ngg lowering send gs_alloc_req at the beginning when culling + * is disabled, but GFX10 may hang if not all waves are launched before gs_alloc_req. + * We work around this HW bug by inserting a barrier before gs_alloc_req. */ - if ((ctx->stage == MESA_SHADER_VERTEX || ctx->stage == MESA_SHADER_TESS_EVAL) && - shader->key.ge.as_ngg && !shader->key.ge.as_es && !shader->key.ge.opt.ngg_culling) { - /* GFX10 requires a barrier before gs_alloc_req due to a hw bug. */ - if (ctx->screen->info.gfx_level == GFX10) - ac_build_s_barrier(&ctx->ac, ctx->stage); - - gfx10_ngg_build_sendmsg_gs_alloc_req(ctx); - - /* Build the primitive export at the beginning - * of the shader if possible. - */ - if (gfx10_ngg_export_prim_early(shader)) - gfx10_ngg_build_export_prim(ctx, NULL, NULL); - } + if (ctx->screen->info.gfx_level == GFX10 && + (ctx->stage == MESA_SHADER_VERTEX || ctx->stage == MESA_SHADER_TESS_EVAL) && + shader->key.ge.as_ngg && !shader->key.ge.as_es && !shader->key.ge.opt.ngg_culling) + ac_build_s_barrier(&ctx->ac, ctx->stage); /* NGG GS: Initialize LDS and insert s_barrier, which must not be inside the if statement. */ if (ctx->stage == MESA_SHADER_GEOMETRY && shader->key.ge.as_ngg) @@ -1164,10 +1157,8 @@ bool si_llvm_translate_nir(struct si_shader_context *ctx, struct si_shader *shad * not here. */ thread_enabled = si_is_gs_thread(ctx); /* 2nd shader: thread enabled bool */ - } else if (((shader->key.ge.as_ls || shader->key.ge.as_es) && !shader->is_monolithic) || - (shader->key.ge.as_ngg && !shader->key.ge.as_es)) { - /* This is NGG VS or NGG TES or VS before GS or TES before GS or VS before TCS. - * For monolithic LS (VS before TCS) and ES (VS before GS and TES before GS), + } else if ((shader->key.ge.as_ls || shader->key.ge.as_es) && !shader->is_monolithic) { + /* For monolithic LS (VS before TCS) and ES (VS before GS and TES before GS), * the if statement is inserted by the wrapper function. */ thread_enabled = si_is_es_thread(ctx); /* 1st shader: thread enabled bool */ @@ -1253,11 +1244,7 @@ bool si_llvm_translate_nir(struct si_shader_context *ctx, struct si_shader *shad si_llvm_ls_build_end(ctx); else if (shader->key.ge.as_es) si_llvm_es_build_end(ctx); - else if (ngg_cull_shader) - gfx10_ngg_culling_build_end(ctx); - else if (shader->key.ge.as_ngg) - gfx10_ngg_build_end(ctx); - else + else if (!shader->key.ge.as_ngg) si_llvm_vs_build_end(ctx); break; @@ -1268,11 +1255,7 @@ bool si_llvm_translate_nir(struct si_shader_context *ctx, struct si_shader *shad case MESA_SHADER_TESS_EVAL: if (ctx->shader->key.ge.as_es) si_llvm_es_build_end(ctx); - else if (ngg_cull_shader) - gfx10_ngg_culling_build_end(ctx); - else if (ctx->shader->key.ge.as_ngg) - gfx10_ngg_build_end(ctx); - else + else if (!ctx->shader->key.ge.as_ngg) si_llvm_vs_build_end(ctx); break; @@ -1323,84 +1306,30 @@ bool si_llvm_compile_shader(struct si_screen *sscreen, struct ac_llvm_compiler * si_llvm_context_init(&ctx, sscreen, compiler, shader->wave_size); ctx.so = *so; - struct ac_llvm_pointer ngg_cull_main_fn = {}; - if (sel->stage <= MESA_SHADER_TESS_EVAL && shader->key.ge.opt.ngg_culling) { - if (!si_llvm_translate_nir(&ctx, shader, nir, false, true)) { - si_llvm_dispose(&ctx); - return false; - } - ngg_cull_main_fn = ctx.main_fn; - ctx.main_fn.value = NULL; - } - - if (!si_llvm_translate_nir(&ctx, shader, nir, free_nir, false)) { + if (!si_llvm_translate_nir(&ctx, shader, nir, free_nir)) { si_llvm_dispose(&ctx); return false; } - if (shader->is_monolithic && sel->stage == MESA_SHADER_VERTEX) { - struct ac_llvm_pointer parts[4]; - unsigned num_parts = 0; - bool first_is_prolog = false; - struct ac_llvm_pointer main_fn = ctx.main_fn; + if (shader->is_monolithic && sel->stage == MESA_SHADER_VERTEX && + si_vs_needs_prolog(sel, &shader->key.ge.part.vs.prolog)) { + struct ac_llvm_pointer parts[2]; + parts[1] = ctx.main_fn; - /* Preserve main arguments. */ - enum ac_arg_type main_arg_types[AC_MAX_ARGS]; - for (int i = 0; i < ctx.args.arg_count; i++) - main_arg_types[i] = ctx.args.args[i].type; - main_arg_types[MIN2(AC_MAX_ARGS - 1, ctx.args.arg_count)] = AC_ARG_INVALID; - - if (ngg_cull_main_fn.value) { - if (si_vs_needs_prolog(sel, &shader->key.ge.part.vs.prolog, &shader->key, true, false)) { - union si_shader_part_key prolog_key; - si_get_vs_prolog_key(&sel->info, shader->info.num_input_sgprs, true, - &shader->key.ge.part.vs.prolog, shader, &prolog_key); - prolog_key.vs_prolog.is_monolithic = true; - si_llvm_build_vs_prolog(&ctx, &prolog_key); - parts[num_parts++] = ctx.main_fn; - first_is_prolog = true; - } - parts[num_parts++] = ngg_cull_main_fn; - } - - if (si_vs_needs_prolog(sel, &shader->key.ge.part.vs.prolog, &shader->key, false, false)) { - union si_shader_part_key prolog_key; - si_get_vs_prolog_key(&sel->info, shader->info.num_input_sgprs, false, - &shader->key.ge.part.vs.prolog, shader, &prolog_key); - prolog_key.vs_prolog.is_monolithic = true; - si_llvm_build_vs_prolog(&ctx, &prolog_key); - parts[num_parts++] = ctx.main_fn; - if (num_parts == 1) - first_is_prolog = true; - } - parts[num_parts++] = main_fn; - - si_build_wrapper_function(&ctx, parts, num_parts, first_is_prolog ? 1 : 0, 0, main_arg_types, false); - } else if (shader->is_monolithic && sel->stage == MESA_SHADER_TESS_EVAL && ngg_cull_main_fn.value) { - struct ac_llvm_pointer parts[3], prolog, main_fn = ctx.main_fn; - - /* Preserve main arguments. */ + /* Preserve main arguments. */ enum ac_arg_type main_arg_types[AC_MAX_ARGS]; for (int i = 0; i < ctx.args.arg_count; i++) main_arg_types[i] = ctx.args.args[i].type; main_arg_types[MIN2(AC_MAX_ARGS - 1, ctx.args.arg_count)] = AC_ARG_INVALID; - /* We reuse the VS prolog code for TES just to load the input VGPRs from LDS. */ union si_shader_part_key prolog_key; - memset(&prolog_key, 0, sizeof(prolog_key)); - prolog_key.vs_prolog.num_input_sgprs = shader->info.num_input_sgprs; - prolog_key.vs_prolog.num_merged_next_stage_vgprs = 5; - prolog_key.vs_prolog.as_ngg = 1; - prolog_key.vs_prolog.load_vgprs_after_culling = 1; + si_get_vs_prolog_key(&sel->info, shader->info.num_input_sgprs, + &shader->key.ge.part.vs.prolog, shader, &prolog_key); prolog_key.vs_prolog.is_monolithic = true; si_llvm_build_vs_prolog(&ctx, &prolog_key); - prolog = ctx.main_fn; - - parts[0] = ngg_cull_main_fn; - parts[1] = prolog; - parts[2] = main_fn; + parts[0] = ctx.main_fn; - si_build_wrapper_function(&ctx, parts, 3, 0, 0, main_arg_types, false); + si_build_wrapper_function(&ctx, parts, 2, 1, 0, main_arg_types, false); } else if (shader->is_monolithic && sel->stage == MESA_SHADER_TESS_CTRL) { /* Preserve main arguments. */ enum ac_arg_type main_arg_types[AC_MAX_ARGS]; @@ -1409,7 +1338,7 @@ bool si_llvm_compile_shader(struct si_screen *sscreen, struct ac_llvm_compiler * struct si_shader_selector *ls = shader->key.ge.part.tcs.ls; struct ac_llvm_pointer parts[4]; bool vs_needs_prolog = - si_vs_needs_prolog(ls, &shader->key.ge.part.tcs.ls_prolog, &shader->key, false, false); + si_vs_needs_prolog(ls, &shader->key.ge.part.tcs.ls_prolog); /* TCS main part */ parts[2] = ctx.main_fn; @@ -1432,7 +1361,7 @@ bool si_llvm_compile_shader(struct si_screen *sscreen, struct ac_llvm_compiler * nir = si_get_nir_shader(&shader_ls, &free_nir, sel->info.tcs_vgpr_only_inputs); si_update_shader_binary_info(shader, nir); - if (!si_llvm_translate_nir(&ctx, &shader_ls, nir, free_nir, false)) { + if (!si_llvm_translate_nir(&ctx, &shader_ls, nir, free_nir)) { si_llvm_dispose(&ctx); return false; } @@ -1446,7 +1375,7 @@ bool si_llvm_compile_shader(struct si_screen *sscreen, struct ac_llvm_compiler * /* LS prolog */ if (vs_needs_prolog) { union si_shader_part_key vs_prolog_key; - si_get_vs_prolog_key(&ls->info, shader_ls.info.num_input_sgprs, false, + si_get_vs_prolog_key(&ls->info, shader_ls.info.num_input_sgprs, &shader->key.ge.part.tcs.ls_prolog, shader, &vs_prolog_key); vs_prolog_key.vs_prolog.is_monolithic = true; si_llvm_build_vs_prolog(&ctx, &vs_prolog_key); @@ -1503,7 +1432,7 @@ bool si_llvm_compile_shader(struct si_screen *sscreen, struct ac_llvm_compiler * nir = si_get_nir_shader(&shader_es, &free_nir, 0); si_update_shader_binary_info(shader, nir); - if (!si_llvm_translate_nir(&ctx, &shader_es, nir, free_nir, false)) { + if (!si_llvm_translate_nir(&ctx, &shader_es, nir, free_nir)) { si_llvm_dispose(&ctx); return false; } @@ -1517,9 +1446,9 @@ bool si_llvm_compile_shader(struct si_screen *sscreen, struct ac_llvm_compiler * /* ES prolog */ if (es->stage == MESA_SHADER_VERTEX && - si_vs_needs_prolog(es, &shader->key.ge.part.gs.vs_prolog, &shader->key, false, true)) { + si_vs_needs_prolog(es, &shader->key.ge.part.gs.vs_prolog)) { union si_shader_part_key vs_prolog_key; - si_get_vs_prolog_key(&es->info, shader_es.info.num_input_sgprs, false, + si_get_vs_prolog_key(&es->info, shader_es.info.num_input_sgprs, &shader->key.ge.part.gs.vs_prolog, shader, &vs_prolog_key); vs_prolog_key.vs_prolog.is_monolithic = true; si_llvm_build_vs_prolog(&ctx, &vs_prolog_key); diff --git a/src/gallium/drivers/radeonsi/si_shader_llvm_gs.c b/src/gallium/drivers/radeonsi/si_shader_llvm_gs.c index c5b1c123231..7deecd389d0 100644 --- a/src/gallium/drivers/radeonsi/si_shader_llvm_gs.c +++ b/src/gallium/drivers/radeonsi/si_shader_llvm_gs.c @@ -462,7 +462,7 @@ struct si_shader *si_generate_gs_copy_shader(struct si_screen *sscreen, builder = ctx.ac.builder; /* Build the main function. */ - si_llvm_create_main_func(&ctx, false); + si_llvm_create_main_func(&ctx); ctx.gsvs_ring[0] = ac_build_load_to_sgpr(&ctx.ac, diff --git a/src/gallium/drivers/radeonsi/si_shader_llvm_vs.c b/src/gallium/drivers/radeonsi/si_shader_llvm_vs.c index 31ba30d5f32..12d45448974 100644 --- a/src/gallium/drivers/radeonsi/si_shader_llvm_vs.c +++ b/src/gallium/drivers/radeonsi/si_shader_llvm_vs.c @@ -698,7 +698,9 @@ void si_llvm_build_vs_exports(struct si_shader_context *ctx, LLVMValueRef num_ex ac_build_export(&ctx->ac, &pos_args[i]); } - if (!shader->info.nr_param_exports) + if (!shader->info.nr_param_exports || + /* GFX11 VS/TES param export is handled in nir */ + (ctx->screen->info.gfx_level >= GFX11 && ctx->stage != MESA_SHADER_GEOMETRY)) return; /* Build parameter exports. Use 2 loops to export params in ascending order. @@ -895,18 +897,6 @@ void si_llvm_build_vs_prolog(struct si_shader_context *ctx, union si_shader_part } } - /* The culling code stored the LDS addresses of the VGPRs into those VGPRs. Load them. */ - if (key->vs_prolog.load_vgprs_after_culling) { - for (i = 5; i <= 8; i++) { - bool is_tes_rel_patch_id = i == 7; - LLVMTypeRef t = is_tes_rel_patch_id ? ctx->ac.i8 : ctx->ac.i32; - input_vgprs[i] = LLVMBuildIntToPtr(ctx->ac.builder, input_vgprs[i], LLVMPointerType(t, AC_ADDR_SPACE_LDS), ""); - input_vgprs[i] = LLVMBuildLoad2(ctx->ac.builder, t, input_vgprs[i], ""); - if (is_tes_rel_patch_id) - input_vgprs[i] = LLVMBuildZExt(ctx->ac.builder, input_vgprs[i], ctx->ac.i32, ""); - } - } - unsigned vertex_id_vgpr = first_vs_vgpr; unsigned instance_id_vgpr = ctx->screen->info.gfx_level >= GFX10 ? first_vs_vgpr + 3 @@ -960,7 +950,7 @@ void si_llvm_build_vs_prolog(struct si_shader_context *ctx, union si_shader_part si_llvm_build_ret(ctx, ret); } -void si_llvm_init_vs_callbacks(struct si_shader_context *ctx, bool ngg_cull_shader) +void si_llvm_init_vs_callbacks(struct si_shader_context *ctx) { ctx->abi.load_inputs = si_load_vs_input; } diff --git a/src/gallium/drivers/radeonsi/si_shader_nir.c b/src/gallium/drivers/radeonsi/si_shader_nir.c index b0c2e4ff5b5..413dce3cfa4 100644 --- a/src/gallium/drivers/radeonsi/si_shader_nir.c +++ b/src/gallium/drivers/radeonsi/si_shader_nir.c @@ -240,6 +240,16 @@ static bool si_lower_intrinsics(nir_shader *nir) NULL); } +const nir_lower_subgroups_options si_nir_subgroups_options = { + .subgroup_size = 64, + .ballot_bit_size = 64, + .ballot_components = 1, + .lower_to_scalar = true, + .lower_subgroup_masks = true, + .lower_vote_trivial = false, + .lower_vote_eq = true, +}; + /** * Perform "lowering" operations on the NIR that are run once when the shader * selector is created. @@ -269,16 +279,7 @@ static void si_lower_nir(struct si_screen *sscreen, struct nir_shader *nir) NIR_PASS_V(nir, si_lower_intrinsics); - const nir_lower_subgroups_options subgroups_options = { - .subgroup_size = 64, - .ballot_bit_size = 64, - .ballot_components = 1, - .lower_to_scalar = true, - .lower_subgroup_masks = true, - .lower_vote_trivial = false, - .lower_vote_eq = true, - }; - NIR_PASS_V(nir, nir_lower_subgroups, &subgroups_options); + NIR_PASS_V(nir, nir_lower_subgroups, &si_nir_subgroups_options); NIR_PASS_V(nir, nir_lower_discard_or_demote, (sscreen->debug_flags & DBG(FS_CORRECT_DERIVS_AFTER_KILL)) ||
