1 /*
2 * Copyright (C) 2020-2023 Collabora, Ltd.
3 *
4 * Permission is hereby granted, free of charge, to any person obtaining a
5 * copy of this software and associated documentation files (the "Software"),
6 * to deal in the Software without restriction, including without limitation
7 * the rights to use, copy, modify, merge, publish, distribute, sublicense,
8 * and/or sell copies of the Software, and to permit persons to whom the
9 * Software is furnished to do so, subject to the following conditions:
10 *
11 * The above copyright notice and this permission notice (including the next
12 * paragraph) shall be included in all copies or substantial portions of the
13 * Software.
14 *
15 * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
16 * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
17 * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL
18 * THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
19 * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
20 * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
21 * SOFTWARE.
22 */
23
24 #include "compiler/nir/nir_builder.h"
25 #include "pan_context.h"
26
27 struct ctx {
28 struct panfrost_sysvals *sysvals;
29 struct hash_table_u64 *sysval_to_id;
30 unsigned sysval_ubo;
31 };
32
33 static unsigned
lookup_sysval(struct hash_table_u64 * sysval_to_id,struct panfrost_sysvals * sysvals,int sysval)34 lookup_sysval(struct hash_table_u64 *sysval_to_id,
35 struct panfrost_sysvals *sysvals, int sysval)
36 {
37 /* Try to lookup */
38 void *cached = _mesa_hash_table_u64_search(sysval_to_id, sysval);
39
40 if (cached) {
41 unsigned id = ((uintptr_t)cached) - 1;
42 assert(id < MAX_SYSVAL_COUNT);
43 assert(sysvals->sysvals[id] == sysval);
44 return id;
45 }
46
47 /* Else assign */
48 unsigned id = sysvals->sysval_count++;
49 assert(id < MAX_SYSVAL_COUNT);
50 _mesa_hash_table_u64_insert(sysval_to_id, sysval,
51 (void *)((uintptr_t)id + 1));
52 sysvals->sysvals[id] = sysval;
53 return id;
54 }
55
56 static unsigned
sysval_for_intrinsic(nir_intrinsic_instr * intr,unsigned * offset)57 sysval_for_intrinsic(nir_intrinsic_instr *intr, unsigned *offset)
58 {
59 switch (intr->intrinsic) {
60 case nir_intrinsic_load_ssbo_address:
61 return PAN_SYSVAL(SSBO, nir_src_as_uint(intr->src[0]));
62 case nir_intrinsic_get_ssbo_size:
63 *offset = 8;
64 return PAN_SYSVAL(SSBO, nir_src_as_uint(intr->src[0]));
65
66 case nir_intrinsic_load_sampler_lod_parameters_pan:
67 /* This is only used for a workaround on Mali-T720, where we don't
68 * support dynamic samplers.
69 */
70 return PAN_SYSVAL(SAMPLER, nir_src_as_uint(intr->src[0]));
71
72 case nir_intrinsic_load_xfb_address:
73 return PAN_SYSVAL(XFB, nir_intrinsic_base(intr));
74
75 case nir_intrinsic_load_work_dim:
76 return PAN_SYSVAL_WORK_DIM;
77
78 case nir_intrinsic_load_sample_positions_pan:
79 return PAN_SYSVAL_SAMPLE_POSITIONS;
80
81 case nir_intrinsic_load_num_vertices:
82 return PAN_SYSVAL_NUM_VERTICES;
83
84 case nir_intrinsic_load_first_vertex:
85 return PAN_SYSVAL_VERTEX_INSTANCE_OFFSETS;
86 case nir_intrinsic_load_base_vertex:
87 *offset = 4;
88 return PAN_SYSVAL_VERTEX_INSTANCE_OFFSETS;
89 case nir_intrinsic_load_base_instance:
90 *offset = 8;
91 return PAN_SYSVAL_VERTEX_INSTANCE_OFFSETS;
92
93 case nir_intrinsic_load_draw_id:
94 return PAN_SYSVAL_DRAWID;
95
96 case nir_intrinsic_load_multisampled_pan:
97 return PAN_SYSVAL_MULTISAMPLED;
98
99 case nir_intrinsic_load_viewport_scale:
100 return PAN_SYSVAL_VIEWPORT_SCALE;
101
102 case nir_intrinsic_load_viewport_offset:
103 return PAN_SYSVAL_VIEWPORT_OFFSET;
104
105 case nir_intrinsic_load_num_workgroups:
106 return PAN_SYSVAL_NUM_WORK_GROUPS;
107
108 case nir_intrinsic_load_workgroup_size:
109 return PAN_SYSVAL_LOCAL_GROUP_SIZE;
110
111 case nir_intrinsic_load_rt_conversion_pan: {
112 unsigned size = nir_alu_type_get_type_size(nir_intrinsic_src_type(intr));
113 unsigned rt = nir_intrinsic_base(intr);
114
115 return PAN_SYSVAL(RT_CONVERSION, rt | (size << 4));
116 }
117
118 case nir_intrinsic_image_size: {
119 uint32_t uindex = nir_src_as_uint(intr->src[0]);
120 bool is_array = nir_intrinsic_image_array(intr);
121 unsigned dim = nir_intrinsic_dest_components(intr) - is_array;
122
123 return PAN_SYSVAL(IMAGE_SIZE, PAN_TXS_SYSVAL_ID(uindex, dim, is_array));
124 }
125
126 default:
127 return ~0;
128 }
129 }
130
131 static bool
lower(nir_builder * b,nir_instr * instr,void * data)132 lower(nir_builder *b, nir_instr *instr, void *data)
133 {
134 struct ctx *ctx = data;
135 nir_def *old = NULL;
136 unsigned sysval = ~0, offset = 0;
137 b->cursor = nir_before_instr(instr);
138
139 if (instr->type == nir_instr_type_intrinsic) {
140 nir_intrinsic_instr *intr = nir_instr_as_intrinsic(instr);
141 old = &intr->def;
142 sysval = sysval_for_intrinsic(intr, &offset);
143
144 if (sysval == ~0)
145 return false;
146 } else if (instr->type == nir_instr_type_tex) {
147 nir_tex_instr *tex = nir_instr_as_tex(instr);
148 old = &tex->def;
149
150 if (tex->op != nir_texop_txs)
151 return false;
152
153 /* XXX: This is broken for dynamic indexing */
154 sysval = PAN_SYSVAL(TEXTURE_SIZE,
155 PAN_TXS_SYSVAL_ID(tex->texture_index,
156 nir_tex_instr_dest_size(tex) -
157 (tex->is_array ? 1 : 0),
158 tex->is_array));
159 } else {
160 return false;
161 }
162
163 /* Allocate a UBO for the sysvals if we haven't yet */
164 if (ctx->sysvals->sysval_count == 0)
165 ctx->sysval_ubo = b->shader->info.num_ubos++;
166
167 unsigned vec4_index = lookup_sysval(ctx->sysval_to_id, ctx->sysvals, sysval);
168 unsigned ubo_offset = (vec4_index * 16) + offset;
169
170 b->cursor = nir_after_instr(instr);
171 nir_def *val = nir_load_ubo(
172 b, old->num_components, old->bit_size, nir_imm_int(b, ctx->sysval_ubo),
173 nir_imm_int(b, ubo_offset), .align_mul = old->bit_size / 8,
174 .align_offset = 0, .range_base = offset, .range = old->bit_size / 8);
175 nir_def_rewrite_uses(old, val);
176 return true;
177 }
178
179 bool
panfrost_nir_lower_sysvals(nir_shader * shader,struct panfrost_sysvals * sysvals)180 panfrost_nir_lower_sysvals(nir_shader *shader, struct panfrost_sysvals *sysvals)
181 {
182 bool progress = false;
183
184 /* The lowerings for SSBOs, etc require constants, so fold now */
185 do {
186 progress = false;
187
188 NIR_PASS(progress, shader, nir_copy_prop);
189 NIR_PASS(progress, shader, nir_opt_constant_folding);
190 NIR_PASS(progress, shader, nir_opt_dce);
191 } while (progress);
192
193 struct ctx ctx = {
194 .sysvals = sysvals,
195 .sysval_to_id = _mesa_hash_table_u64_create(NULL),
196 };
197
198 memset(sysvals, 0, sizeof(*sysvals));
199
200 nir_shader_instructions_pass(
201 shader, lower, nir_metadata_block_index | nir_metadata_dominance, &ctx);
202
203 _mesa_hash_table_u64_destroy(ctx.sysval_to_id);
204 return true;
205 }
206