1 /* -*- mesa-c++ -*-
2 * Copyright 2022 Collabora LTD
3 * Author: Gert Wollny <[email protected]>
4 * SPDX-License-Identifier: MIT
5 */
6
7 #include "sfn_shader_cs.h"
8
9 #include "sfn_instr_fetch.h"
10
11 namespace r600 {
12
ComputeShader(UNUSED const r600_shader_key & key,int num_samplers)13 ComputeShader::ComputeShader(UNUSED const r600_shader_key& key, int num_samplers):
14 Shader("CS", 0),
15 m_image_size_const_offset(num_samplers)
16 {
17 }
18
19 bool
do_scan_instruction(UNUSED nir_instr * instr)20 ComputeShader::do_scan_instruction(UNUSED nir_instr *instr)
21 {
22 return false;
23 }
24
25 int
do_allocate_reserved_registers()26 ComputeShader::do_allocate_reserved_registers()
27 {
28 auto& vf = value_factory();
29
30 const int thread_id_sel = 0;
31 const int wg_id_sel = 1;
32
33 for (int i = 0; i < 3; ++i) {
34 m_local_invocation_id[i] = vf.allocate_pinned_register(thread_id_sel, i);
35 m_local_invocation_id[i]->set_flag(Register::pin_end);
36 m_workgroup_id[i] = vf.allocate_pinned_register(wg_id_sel, i);
37 m_workgroup_id[i]->set_flag(Register::pin_end);
38 }
39 return 2;
40 }
41
42 bool
process_stage_intrinsic(nir_intrinsic_instr * instr)43 ComputeShader::process_stage_intrinsic(nir_intrinsic_instr *instr)
44 {
45 switch (instr->intrinsic) {
46 case nir_intrinsic_load_local_invocation_id:
47 return emit_load_3vec(instr, m_local_invocation_id);
48 case nir_intrinsic_load_workgroup_id:
49 return emit_load_3vec(instr, m_workgroup_id);
50 case nir_intrinsic_load_workgroup_size:
51 return emit_load_from_info_buffer(instr, 0);
52 case nir_intrinsic_load_num_workgroups:
53 return emit_load_from_info_buffer(instr, 16);
54 default:
55 return false;
56 }
57 }
58
59 void
do_get_shader_info(r600_shader * sh_info)60 ComputeShader::do_get_shader_info(r600_shader *sh_info)
61 {
62 sh_info->processor_type = PIPE_SHADER_COMPUTE;
63 }
64
65 bool
read_prop(UNUSED std::istream & is)66 ComputeShader::read_prop(UNUSED std::istream& is)
67 {
68 return true;
69 }
70
71 void
do_print_properties(UNUSED std::ostream & os) const72 ComputeShader::do_print_properties(UNUSED std::ostream& os) const
73 {
74 }
75
76 bool
emit_load_from_info_buffer(nir_intrinsic_instr * instr,int offset)77 ComputeShader::emit_load_from_info_buffer(nir_intrinsic_instr *instr, int offset)
78 {
79 if (!m_zero_register) {
80 m_zero_register = value_factory().temp_register();
81 emit_instruction(new AluInstr(op1_mov,
82 m_zero_register,
83 value_factory().inline_const(ALU_SRC_0, 0),
84 AluInstr::last_write));
85 }
86
87 auto dest = value_factory().dest_vec4(instr->def, pin_group);
88
89 auto ir = new LoadFromBuffer(dest,
90 {0, 1, 2, 7},
91 m_zero_register,
92 offset,
93 R600_BUFFER_INFO_CONST_BUFFER,
94 nullptr,
95 fmt_32_32_32_32);
96
97 ir->set_fetch_flag(LoadFromBuffer::srf_mode);
98 ir->reset_fetch_flag(LoadFromBuffer::format_comp_signed);
99 ir->set_num_format(vtx_nf_int);
100 emit_instruction(ir);
101 return true;
102 }
103
104 bool
emit_load_3vec(nir_intrinsic_instr * instr,const std::array<PRegister,3> & src)105 ComputeShader::emit_load_3vec(nir_intrinsic_instr *instr,
106 const std::array<PRegister, 3>& src)
107 {
108 auto& vf = value_factory();
109
110 for (int i = 0; i < 3; ++i) {
111 auto dest = vf.dest(instr->def, i, pin_none);
112 emit_instruction(new AluInstr(
113 op1_mov, dest, src[i], i == 2 ? AluInstr::last_write : AluInstr::write));
114 }
115 return true;
116 }
117
118 } // namespace r600
119