Module: Mesa Branch: main Commit: 3b9f36db47702a49e43beaeabe7025d73503a2a3 URL: http://cgit.freedesktop.org/mesa/mesa/commit/?id=3b9f36db47702a49e43beaeabe7025d73503a2a3
Author: Gert Wollny <[email protected]> Date: Sun Oct 30 10:20:13 2022 +0100 r600/sfn: Handle load_workgroup_size Fixes: 79ca456b4837b3bc21cf9ef3c03c505c4b4909f6 r600/sfn: rewrite NIR backend Signed-off-by: Gert Wollny <[email protected]> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/19417> --- src/gallium/drivers/r600/sfn/sfn_shader_cs.cpp | 20 +++++++++++++------- src/gallium/drivers/r600/sfn/sfn_shader_cs.h | 4 +++- 2 files changed, 16 insertions(+), 8 deletions(-) diff --git a/src/gallium/drivers/r600/sfn/sfn_shader_cs.cpp b/src/gallium/drivers/r600/sfn/sfn_shader_cs.cpp index ad81f35c7c0..2ebf976fa54 100644 --- a/src/gallium/drivers/r600/sfn/sfn_shader_cs.cpp +++ b/src/gallium/drivers/r600/sfn/sfn_shader_cs.cpp @@ -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); diff --git a/src/gallium/drivers/r600/sfn/sfn_shader_cs.h b/src/gallium/drivers/r600/sfn/sfn_shader_cs.h index ae4bb6031a0..49bb3211446 100644 --- a/src/gallium/drivers/r600/sfn/sfn_shader_cs.h +++ b/src/gallium/drivers/r600/sfn/sfn_shader_cs.h @@ -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
