diff --git a/src/amd/common/ac_nir.c b/src/amd/common/ac_nir.c index 59f766f401d..5814175202c 100644 --- a/src/amd/common/ac_nir.c +++ b/src/amd/common/ac_nir.c @@ -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; } diff --git a/src/amd/compiler/aco_instruction_selection.cpp b/src/amd/compiler/aco_instruction_selection.cpp index 54c6012b623..56d55e16a00 100644 --- a/src/amd/compiler/aco_instruction_selection.cpp +++ b/src/amd/compiler/aco_instruction_selection.cpp @@ -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); } diff --git a/src/amd/compiler/aco_instruction_selection_setup.cpp b/src/amd/compiler/aco_instruction_selection_setup.cpp index f30461ec0b9..14eb1448ac4 100644 --- a/src/amd/compiler/aco_instruction_selection_setup.cpp +++ b/src/amd/compiler/aco_instruction_selection_setup.cpp @@ -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: diff --git a/src/amd/llvm/ac_nir_to_llvm.c b/src/amd/llvm/ac_nir_to_llvm.c index c5ae7c9d531..570025121b9 100644 --- a/src/amd/llvm/ac_nir_to_llvm.c +++ b/src/amd/llvm/ac_nir_to_llvm.c @@ -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, ""); -} diff --git a/src/amd/llvm/ac_nir_to_llvm.h b/src/amd/llvm/ac_nir_to_llvm.h index 43a69f99acd..a9598319e0d 100644 --- a/src/amd/llvm/ac_nir_to_llvm.h +++ b/src/amd/llvm/ac_nir_to_llvm.h @@ -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 */ diff --git a/src/amd/llvm/ac_shader_abi.h b/src/amd/llvm/ac_shader_abi.h index b2c16717960..7e08007d663 100644 --- a/src/amd/llvm/ac_shader_abi.h +++ b/src/amd/llvm/ac_shader_abi.h @@ -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; diff --git a/src/amd/vulkan/radv_nir_to_llvm.c b/src/amd/vulkan/radv_nir_to_llvm.c index 1fc854388ea..226c84e62da 100644 --- a/src/amd/vulkan/radv_nir_to_llvm.c +++ b/src/amd/vulkan/radv_nir_to_llvm.c @@ -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); diff --git a/src/gallium/drivers/radeonsi/si_shader_llvm.c b/src/gallium/drivers/radeonsi/si_shader_llvm.c index ae78628f07a..53f3d16ba71 100644 --- a/src/gallium/drivers/radeonsi/si_shader_llvm.c +++ b/src/gallium/drivers/radeonsi/si_shader_llvm.c @@ -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)