fdo-mirrors/mesa
1
0
Fork 0
mirror of https://gitlab.freedesktop.org/mesa/mesa.git synced 2026-05-05 11:48:06 +02:00
Code Issues Projects Releases Packages Wiki Activity Actions 2

r600/sfn: Handle load_workgroup_size

Browse source 
Fixes: 79ca456b48
   r600/sfn: rewrite NIR backend

Signed-off-by: Gert Wollny <gert.wollny@collabora.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/19417>
This commit is contained in:
Gert Wollny 2022-10-30 10:20:13 +01:00 • committed by Marge Bot
parent 76555a4777
commit 3b9f36db47
2 changed files with 16 additions and 8 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

20
src/gallium/drivers/r600/sfn/sfn_shader_cs.cpp
View file

@ -67,8 +67,10 @@ ComputeShader::process_stage_intrinsic(nir_intrinsic_instr *instr)
return emit_load_3vec(instr, m_local_invocation_id);
case nir_intrinsic_load_workgroup_id:
return emit_load_3vec(instr, m_workgroup_id);
case nir_intrinsic_load_workgroup_size:
return emit_load_from_info_buffer(instr, 0);
case nir_intrinsic_load_num_workgroups:
return emit_load_num_workgroups(instr);
return emit_load_from_info_buffer(instr, 16);
default:
return false;
}
@ -92,18 +94,22 @@ ComputeShader::do_print_properties(UNUSED std::ostream& os) const
}
bool
ComputeShader::emit_load_num_workgroups(nir_intrinsic_instr *instr)
ComputeShader::emit_load_from_info_buffer(nir_intrinsic_instr *instr, int offset)
{
auto zero = value_factory().temp_register();
if (!m_zero_register) {
m_zero_register = value_factory().temp_register();
emit_instruction(new AluInstr(op1_mov,
m_zero_register,
value_factory().inline_const(ALU_SRC_0, 0),
AluInstr::last_write));
}
emit_instruction(new AluInstr(
op1_mov, zero, value_factory().inline_const(ALU_SRC_0, 0), AluInstr::last_write));
auto dest = value_factory().dest_vec4(instr->dest, pin_group);
auto ir = new LoadFromBuffer(dest,
{0, 1, 2, 7},
zero,
16,
m_zero_register,
offset,
R600_BUFFER_INFO_CONST_BUFFER,
nullptr,
fmt_32_32_32_32);

4
src/gallium/drivers/r600/sfn/sfn_shader_cs.h
View file

@ -54,11 +54,13 @@ private:
bool read_prop(std::istream& is) override;
void do_print_properties(std::ostream& os) const override;
bool emit_load_num_workgroups(nir_intrinsic_instr *instr);
bool emit_load_from_info_buffer(nir_intrinsic_instr *instr, int offset);
bool emit_load_3vec(nir_intrinsic_instr *instr, const std::array<PRegister, 3>& src);
std::array<PRegister, 3> m_workgroup_id{nullptr};
std::array<PRegister, 3> m_local_invocation_id{nullptr};
PRegister m_zero_register{0};
};
} // namespace r600