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

amd: lower load_local_invocation_index in NIR

Browse Source 
This is the last intrinsic that needed the LS VGPR bug workaround in ACO
and ac_nir_to_llvm.

Reviewed-by: Timur Kristóf <timur.kristof@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/32782>
This commit is contained in:
Marek Olšák
2024-12-25 12:44:59 -05:00
committed by Marge Bot
parent 433ca6ba38
commit e69f47faee
8 changed files with 24 additions and 150 deletions
Download Patch File Download Diff File Expand all files Collapse all files
src/amd/common/ac_nir.c
+24
View File
@@ -105,6 +105,7 @@ typedef struct {
nir_def *vertex_id;
nir_def *instance_id;
nir_def *vs_rel_patch_id;
} lower_intrinsics_to_args_state;
static nir_def *
@@ -377,6 +378,29 @@ lower_intrinsic_to_arg(nir_builder *b, nir_instr *instr, void *state)
s->instance_id = preload_arg(s, b->impl, s->args->instance_id, s->args->vertex_id);
replacement = s->instance_id;
break;
case nir_intrinsic_load_local_invocation_index:
/* GFX11 HS has subgroup_id, so use it instead of vs_rel_patch_id. */
if (s->gfx_level < GFX11 &&
(s->hw_stage == AC_HW_LOCAL_SHADER || s->hw_stage == AC_HW_HULL_SHADER)) {
if (!s->vs_rel_patch_id) {
s->vs_rel_patch_id = preload_arg(s, b->impl, s->args->vs_rel_patch_id,
s->args->tcs_rel_ids);
}
replacement = s->vs_rel_patch_id;
} else if (s->workgroup_size <= s->wave_size) {
/* Just a subgroup invocation ID. */
replacement = nir_mbcnt_amd(b, nir_imm_intN_t(b, ~0ull, s->wave_size), nir_imm_int(b, 0));
} else if (s->gfx_level < GFX12 && s->hw_stage == AC_HW_COMPUTE_SHADER && s->wave_size == 64) {
/* After the AND the bits are already multiplied by 64 (left shifted by 6) so we can just
* feed that to mbcnt. (GFX12 doesn't have tg_size)
*/
nir_def *wave_id_mul_64 = nir_iand_imm(b, ac_nir_load_arg(b, s->args, s->args->tg_size), 0xfc0);
replacement = nir_mbcnt_amd(b, nir_imm_intN_t(b, ~0ull, s->wave_size), wave_id_mul_64);
} else {
replacement = nir_mbcnt_amd(b, nir_imm_intN_t(b, ~0ull, s->wave_size),
nir_imul_imm(b, load_subgroup_id_lowered(s, b), s->wave_size));
}
break;
default:
return false;
}
src/amd/compiler/aco_instruction_selection.cpp
-103
View File
@@ -5277,32 +5277,6 @@ resolve_excess_vmem_const_offset(Builder& bld, Temp& voffset, unsigned const_off
return const_offset;
}
Temp
wave_id_in_threadgroup(isel_context* ctx)
{
Builder bld(ctx->program, ctx->block);
return bld.sop2(aco_opcode::s_bfe_u32, bld.def(s1), bld.def(s1, scc),
get_arg(ctx, ctx->args->merged_wave_info), Operand::c32(24u | (4u << 16)));
}
Temp
thread_id_in_threadgroup(isel_context* ctx)
{
/* tid_in_tg = wave_id * wave_size + tid_in_wave */
Builder bld(ctx->program, ctx->block);
Temp tid_in_wave = emit_mbcnt(ctx, bld.tmp(v1));
if (ctx->program->workgroup_size <= ctx->program->wave_size)
return tid_in_wave;
Temp wave_id_in_tg = wave_id_in_threadgroup(ctx);
Temp num_pre_threads =
bld.sop2(aco_opcode::s_lshl_b32, bld.def(s1), bld.def(s1, scc), wave_id_in_tg,
Operand::c32(ctx->program->wave_size == 64 ? 6u : 5u));
return bld.vadd32(bld.def(v1), Operand(num_pre_threads), Operand(tid_in_wave));
}
bool
store_output_to_temps(isel_context* ctx, nir_intrinsic_instr* instr)
{
@@ -8070,60 +8044,6 @@ visit_intrinsic(isel_context* ctx, nir_intrinsic_instr* instr)
ctx->ttmp8, Operand::c32(25 | (5 << 16)));
break;
}
case nir_intrinsic_load_local_invocation_index: {
if (ctx->stage.hw == AC_HW_LOCAL_SHADER || ctx->stage.hw == AC_HW_HULL_SHADER) {
if (ctx->options->gfx_level >= GFX11) {
/* On GFX11, RelAutoIndex is WaveID * WaveSize + ThreadID. */
Temp wave_id =
bld.sop2(aco_opcode::s_bfe_u32, bld.def(s1), bld.def(s1, scc),
get_arg(ctx, ctx->args->tcs_wave_id), Operand::c32(0u | (3u << 16)));
Temp temp = bld.sop2(aco_opcode::s_mul_i32, bld.def(s1), wave_id,
Operand::c32(ctx->program->wave_size));
emit_mbcnt(ctx, get_ssa_temp(ctx, &instr->def), Operand(), Operand(temp));
} else {
bld.copy(Definition(get_ssa_temp(ctx, &instr->def)),
get_arg(ctx, ctx->args->vs_rel_patch_id));
}
break;
} else if (ctx->stage.hw == AC_HW_LEGACY_GEOMETRY_SHADER ||
ctx->stage.hw == AC_HW_NEXT_GEN_GEOMETRY_SHADER) {
bld.copy(Definition(get_ssa_temp(ctx, &instr->def)), thread_id_in_threadgroup(ctx));
break;
} else if (ctx->program->workgroup_size <= ctx->program->wave_size) {
emit_mbcnt(ctx, get_ssa_temp(ctx, &instr->def));
break;
}
Temp id = emit_mbcnt(ctx, bld.tmp(v1));
if (ctx->options->gfx_level >= GFX12) {
Temp tg_num = bld.sop2(aco_opcode::s_bfe_u32, bld.def(s1), bld.def(s1, scc), ctx->ttmp8,
Operand::c32(25 | (5 << 16)));
bld.vop3(aco_opcode::v_lshl_or_b32, Definition(get_ssa_temp(ctx, &instr->def)), tg_num,
Operand::c32(ctx->program->wave_size == 64 ? 6 : 5), id);
break;
}
/* The tg_size bits [6:11] contain the subgroup id,
* we need this multiplied by the wave size, and then OR the thread id to it.
*/
if (ctx->program->wave_size == 64) {
/* After the s_and the bits are already multiplied by 64 (left shifted by 6) so we can just
* feed that to v_or */
Temp tg_num = bld.sop2(aco_opcode::s_and_b32, bld.def(s1), bld.def(s1, scc),
Operand::c32(0xfc0u), get_arg(ctx, ctx->args->tg_size));
bld.vop2(aco_opcode::v_or_b32, Definition(get_ssa_temp(ctx, &instr->def)), tg_num, id);
} else {
/* Extract the bit field and multiply the result by 32 (left shift by 5), then do the OR */
Temp tg_num =
bld.sop2(aco_opcode::s_bfe_u32, bld.def(s1), bld.def(s1, scc),
get_arg(ctx, ctx->args->tg_size), Operand::c32(0x6u | (0x6u << 16)));
bld.vop3(aco_opcode::v_lshl_or_b32, Definition(get_ssa_temp(ctx, &instr->def)), tg_num,
Operand::c32(0x5u), id);
}
break;
}
case nir_intrinsic_ddx:
case nir_intrinsic_ddy:
case nir_intrinsic_ddx_fine:
@@ -11092,24 +11012,6 @@ add_startpgm(struct isel_context* ctx)
return startpgm;
}
void
fix_ls_vgpr_init_bug(isel_context* ctx)
{
Builder bld(ctx->program, ctx->block);
constexpr unsigned hs_idx = 1u;
Builder::Result hs_thread_count =
bld.sop2(aco_opcode::s_bfe_u32, bld.def(s1), bld.def(s1, scc),
get_arg(ctx, ctx->args->merged_wave_info), Operand::c32((8u << 16) | (hs_idx * 8u)));
Temp ls_has_nonzero_hs_threads = bool_to_vector_condition(ctx, hs_thread_count.def(1).getTemp());
/* If there are no HS threads, SPI mistakenly loads the LS VGPRs starting at VGPR 0. */
Temp vs_rel_patch_id =
bld.vop2(aco_opcode::v_cndmask_b32, bld.def(v1), get_arg(ctx, ctx->args->tcs_rel_ids),
get_arg(ctx, ctx->args->vs_rel_patch_id), ls_has_nonzero_hs_threads);
ctx->arg_temps[ctx->args->vs_rel_patch_id.arg_index] = vs_rel_patch_id;
}
void
split_arguments(isel_context* ctx, Instruction* startpgm)
{
@@ -11535,11 +11437,6 @@ select_shader(isel_context& ctx, nir_shader* nir, const bool need_startpgm, cons
}
append_logical_start(ctx.block);
if (ctx.options->has_ls_vgpr_init_bug && ctx.stage == vertex_tess_control_hs &&
!program->info.vs.has_prolog)
fix_ls_vgpr_init_bug(&ctx);
split_arguments(&ctx, startpgm);
}
src/amd/compiler/aco_instruction_selection_setup.cpp
-1
View File
@@ -550,7 +550,6 @@ init_context(isel_context* ctx, nir_shader* shader)
case nir_intrinsic_load_per_vertex_input:
case nir_intrinsic_load_per_vertex_output:
case nir_intrinsic_load_interpolated_input:
case nir_intrinsic_load_local_invocation_index:
case nir_intrinsic_load_subgroup_invocation:
case nir_intrinsic_load_tess_coord:
case nir_intrinsic_write_invocation_amd:
src/amd/llvm/ac_nir_to_llvm.c
-25
View File
@@ -2506,16 +2506,6 @@ static LLVMValueRef visit_load_subgroup_id(struct ac_nir_context *ctx)
}
}
static LLVMValueRef visit_load_local_invocation_index(struct ac_nir_context *ctx)
{
if (ctx->abi->vs_rel_patch_id)
return ctx->abi->vs_rel_patch_id;
return ac_build_imad(&ctx->ac, visit_load_subgroup_id(ctx),
LLVMConstInt(ctx->ac.i32, ctx->ac.wave_size, 0),
ac_get_thread_id(&ctx->ac));
}
static LLVMValueRef visit_first_invocation(struct ac_nir_context *ctx)
{
LLVMValueRef active_set = ac_build_ballot(&ctx->ac, ctx->ac.i32_1);
@@ -2853,9 +2843,6 @@ static bool visit_intrinsic(struct ac_nir_context *ctx, nir_intrinsic_instr *ins
case nir_intrinsic_load_subgroup_id:
result = visit_load_subgroup_id(ctx);
break;
case nir_intrinsic_load_local_invocation_index:
result = visit_load_local_invocation_index(ctx);
break;
case nir_intrinsic_first_invocation:
result = visit_first_invocation(ctx);
break;
@@ -4188,15 +4175,3 @@ bool ac_nir_translate(struct ac_llvm_context *ac, struct ac_shader_abi *abi,
return ret;
}
/* Fixup the HW not emitting the TCS regs if there are no HS threads. */
void ac_fixup_ls_hs_input_vgprs(struct ac_llvm_context *ac, struct ac_shader_abi *abi,
const struct ac_shader_args *args)
{
LLVMValueRef count = ac_unpack_param(ac, ac_get_arg(ac, args->merged_wave_info), 8, 8);
LLVMValueRef hs_empty = LLVMBuildICmp(ac->builder, LLVMIntEQ, count, ac->i32_0, "");
abi->vs_rel_patch_id =
LLVMBuildSelect(ac->builder, hs_empty, ac_get_arg(ac, args->tcs_rel_ids),
abi->vs_rel_patch_id, "");
}
src/amd/llvm/ac_nir_to_llvm.h
-3
View File
@@ -27,7 +27,4 @@ static inline unsigned ac_llvm_reg_index_soa(unsigned index, unsigned chan)
bool ac_nir_translate(struct ac_llvm_context *ac, struct ac_shader_abi *abi,
const struct ac_shader_args *args, struct nir_shader *nir);
void ac_fixup_ls_hs_input_vgprs(struct ac_llvm_context *ac, struct ac_shader_abi *abi,
const struct ac_shader_args *args);
#endif /* AC_NIR_TO_LLVM_H */
src/amd/llvm/ac_shader_abi.h
-3
View File
@@ -25,9 +25,6 @@ struct ac_shader_abi {
LLVMValueRef outputs[AC_LLVM_MAX_OUTPUTS * 4];
bool is_16bit[AC_LLVM_MAX_OUTPUTS * 4];
/* These input registers sometimes need to be fixed up. */
LLVMValueRef vs_rel_patch_id;
/* replaced registers when culling enabled */
LLVMValueRef tes_u_replaced;
LLVMValueRef tes_v_replaced;
src/amd/vulkan/radv_nir_to_llvm.c
-6
View File
@@ -279,12 +279,6 @@ ac_translate_nir_to_llvm(struct ac_llvm_compiler *ac_llvm, const struct radv_nir
if (shader_count >= 2 || is_ngg)
ac_init_exec_full_mask(&ctx.ac);
if (args->ac.vs_rel_patch_id.used)
ctx.abi.vs_rel_patch_id = ac_get_arg(&ctx.ac, args->ac.vs_rel_patch_id);
if (options->info->has_ls_vgpr_init_bug && shaders[shader_count - 1]->info.stage == MESA_SHADER_TESS_CTRL)
ac_fixup_ls_hs_input_vgprs(&ctx.ac, &ctx.abi, &args->ac);
if (is_ngg) {
if (!info->is_ngg_passthrough)
declare_esgs_ring(&ctx);
src/gallium/drivers/radeonsi/si_shader_llvm.c
-9
View File
@@ -217,15 +217,6 @@ void si_llvm_create_main_func(struct si_shader_context *ctx)
};
LLVMSetAlignment(ctx->ac.lds.value, 256);
}
if (ctx->stage == MESA_SHADER_VERTEX) {
if (ctx->args->ac.vs_rel_patch_id.used)
ctx->abi.vs_rel_patch_id = ac_get_arg(&ctx->ac, ctx->args->ac.vs_rel_patch_id);
/* Apply the LS-HS input VGPR hw bug workaround. */
if (shader->key.ge.as_ls && ctx->screen->info.has_ls_vgpr_init_bug)
ac_fixup_ls_hs_input_vgprs(&ctx->ac, &ctx->abi, &ctx->args->ac);
}
}
void si_llvm_optimize_module(struct si_shader_context *ctx)