1 /*
2 * Copyright 2022 Alyssa Rosenzweig
3 * Copyright 2021 Valve Corporation
4 * SPDX-License-Identifier: MIT
5 */
6
7 #include "compiler/nir/nir_builder.h"
8 #include "agx_compiler.h"
9
10 static void
def_size(nir_def * def,unsigned * size,unsigned * align)11 def_size(nir_def *def, unsigned *size, unsigned *align)
12 {
13 unsigned bit_size = MAX2(def->bit_size, 16);
14
15 *size = (bit_size * def->num_components) / 16;
16 *align = bit_size / 16;
17 }
18
19 static float
instr_cost(nir_instr * instr,const void * data)20 instr_cost(nir_instr *instr, const void *data)
21 {
22 switch (instr->type) {
23 case nir_instr_type_intrinsic:
24 switch (nir_instr_as_intrinsic(instr)->intrinsic) {
25 case nir_intrinsic_load_global:
26 case nir_intrinsic_load_agx:
27 case nir_intrinsic_load_global_constant:
28 case nir_intrinsic_load_constant_agx:
29 case nir_intrinsic_load_ubo:
30 return 10.0;
31 default:
32 /* Assume it's a sysval or something */
33 return 0.0;
34 }
35
36 case nir_instr_type_tex:
37 /* Texturing involes lots of memory bandwidth */
38 return 20.0;
39
40 case nir_instr_type_alu:
41 /* We optimistically assume that moves get coalesced */
42 if (nir_op_is_vec_or_mov(nir_instr_as_alu(instr)->op))
43 return 0.0;
44 else
45 return 2.0;
46
47 default:
48 return 1.0;
49 }
50 }
51
52 static float
rewrite_cost(nir_def * def,const void * data)53 rewrite_cost(nir_def *def, const void *data)
54 {
55 bool mov_needed = false;
56 nir_foreach_use(use, def) {
57 nir_instr *parent_instr = nir_src_parent_instr(use);
58 if (parent_instr->type != nir_instr_type_alu) {
59 mov_needed = true;
60 break;
61 } else {
62 nir_alu_instr *alu = nir_instr_as_alu(parent_instr);
63 if (alu->op == nir_op_vec2 || alu->op == nir_op_vec3 ||
64 alu->op == nir_op_vec4 || alu->op == nir_op_mov) {
65 mov_needed = true;
66 break;
67 } else {
68 /* Assume for non-moves that the const is folded into the src */
69 }
70 }
71 }
72
73 return mov_needed ? def->num_components : 0;
74 }
75
76 static bool
avoid_instr(const nir_instr * instr,const void * data)77 avoid_instr(const nir_instr *instr, const void *data)
78 {
79 const nir_def *def = nir_instr_def((nir_instr *)instr);
80
81 /* Do not move bindless handles, since we need those to retain their constant
82 * base index.
83 */
84 if (def) {
85 nir_foreach_use(use, def) {
86 if (nir_src_parent_instr(use)->type == nir_instr_type_tex) {
87 /* Check if used as a bindless texture handle */
88 nir_tex_instr *tex = nir_instr_as_tex(nir_src_parent_instr(use));
89 int handle_idx =
90 nir_tex_instr_src_index(tex, nir_tex_src_texture_handle);
91
92 if (handle_idx >= 0 && tex->src[handle_idx].src.ssa == def)
93 return true;
94 } else if (nir_src_parent_instr(use)->type ==
95 nir_instr_type_intrinsic) {
96 /* Check if used as a bindless image handle */
97 nir_intrinsic_instr *intr =
98 nir_instr_as_intrinsic(nir_src_parent_instr(use));
99
100 switch (intr->intrinsic) {
101 case nir_intrinsic_bindless_image_load:
102 case nir_intrinsic_bindless_image_store:
103 if (intr->src[0].ssa == def)
104 return true;
105 break;
106 default:
107 break;
108 }
109 }
110 }
111 }
112
113 return false;
114 }
115
116 static const nir_opt_preamble_options preamble_options = {
117 .drawid_uniform = true,
118 .subgroup_size_uniform = true,
119 /* not supported in hardware */
120 .load_workgroup_size_allowed = false,
121 .def_size = def_size,
122 .instr_cost_cb = instr_cost,
123 .rewrite_cost_cb = rewrite_cost,
124 .avoid_instr_cb = avoid_instr,
125 .preamble_storage_size = 512,
126 };
127
128 bool
agx_nir_opt_preamble(nir_shader * nir,unsigned * preamble_size)129 agx_nir_opt_preamble(nir_shader *nir, unsigned *preamble_size)
130 {
131 return nir_opt_preamble(nir, &preamble_options, preamble_size);
132 }
133