projects / mesa.git / blob
? search:
summary | shortlog | log | commit | commitdiff | tree
history | raw | HEAD
radeonsi: add a common function for getting the size of gs_ngg_scratch
[mesa.git] / src / gallium / drivers / radeonsi / si_shader.c
1 /*
2 * Copyright 2012 Advanced Micro Devices, Inc.
3 * All Rights Reserved.
4 *
5 * Permission is hereby granted, free of charge, to any person obtaining a
6 * copy of this software and associated documentation files (the "Software"),
7 * to deal in the Software without restriction, including without limitation
8 * on the rights to use, copy, modify, merge, publish, distribute, sub
9 * license, and/or sell copies of the Software, and to permit persons to whom
10 * the Software is furnished to do so, subject to the following conditions:
11 *
12 * The above copyright notice and this permission notice (including the next
13 * paragraph) shall be included in all copies or substantial portions of the
14 * Software.
15 *
16 * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
17 * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
18 * FITNESS FOR A PARTICULAR PURPOSE AND NON-INFRINGEMENT. IN NO EVENT SHALL
19 * THE AUTHOR(S) AND/OR THEIR SUPPLIERS BE LIABLE FOR ANY CLAIM,
20 * DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR
21 * OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE
22 * USE OR OTHER DEALINGS IN THE SOFTWARE.
23 */
24
25 #include "ac_exp_param.h"
26 #include "ac_rtld.h"
27 #include "compiler/nir/nir.h"
28 #include "compiler/nir/nir_serialize.h"
29 #include "si_pipe.h"
30 #include "si_shader_internal.h"
31 #include "sid.h"
32 #include "tgsi/tgsi_from_mesa.h"
33 #include "tgsi/tgsi_strings.h"
34 #include "util/u_memory.h"
35
36 static const char scratch_rsrc_dword0_symbol[] = "SCRATCH_RSRC_DWORD0";
37
38 static const char scratch_rsrc_dword1_symbol[] = "SCRATCH_RSRC_DWORD1";
39
40 static void si_dump_shader_key(const struct si_shader *shader, FILE *f);
41
42 /** Whether the shader runs as a combination of multiple API shaders */
43 bool si_is_multi_part_shader(struct si_shader *shader)
44 {
45 if (shader->selector->screen->info.chip_class <= GFX8)
46 return false;
47
48 return shader->key.as_ls || shader->key.as_es ||
49 shader->selector->type == PIPE_SHADER_TESS_CTRL ||
50 shader->selector->type == PIPE_SHADER_GEOMETRY;
51 }
52
53 /** Whether the shader runs on a merged HW stage (LSHS or ESGS) */
54 bool si_is_merged_shader(struct si_shader *shader)
55 {
56 return shader->key.as_ngg || si_is_multi_part_shader(shader);
57 }
58
59 /**
60 * Returns a unique index for a per-patch semantic name and index. The index
61 * must be less than 32, so that a 32-bit bitmask of used inputs or outputs
62 * can be calculated.
63 */
64 unsigned si_shader_io_get_unique_index_patch(unsigned semantic_name, unsigned index)
65 {
66 switch (semantic_name) {
67 case TGSI_SEMANTIC_TESSOUTER:
68 return 0;
69 case TGSI_SEMANTIC_TESSINNER:
70 return 1;
71 case TGSI_SEMANTIC_PATCH:
72 assert(index < 30);
73 return 2 + index;
74
75 default:
76 assert(!"invalid semantic name");
77 return 0;
78 }
79 }
80
81 /**
82 * Returns a unique index for a semantic name and index. The index must be
83 * less than 64, so that a 64-bit bitmask of used inputs or outputs can be
84 * calculated.
85 */
86 unsigned si_shader_io_get_unique_index(unsigned semantic_name, unsigned index, unsigned is_varying)
87 {
88 switch (semantic_name) {
89 case TGSI_SEMANTIC_POSITION:
90 return 0;
91 case TGSI_SEMANTIC_GENERIC:
92 /* Since some shader stages use the the highest used IO index
93 * to determine the size to allocate for inputs/outputs
94 * (in LDS, tess and GS rings). GENERIC should be placed right
95 * after POSITION to make that size as small as possible.
96 */
97 if (index < SI_MAX_IO_GENERIC)
98 return 1 + index;
99
100 assert(!"invalid generic index");
101 return 0;
102 case TGSI_SEMANTIC_FOG:
103 return SI_MAX_IO_GENERIC + 1;
104 case TGSI_SEMANTIC_COLOR:
105 assert(index < 2);
106 return SI_MAX_IO_GENERIC + 2 + index;
107 case TGSI_SEMANTIC_BCOLOR:
108 assert(index < 2);
109 /* If it's a varying, COLOR and BCOLOR alias. */
110 if (is_varying)
111 return SI_MAX_IO_GENERIC + 2 + index;
112 else
113 return SI_MAX_IO_GENERIC + 4 + index;
114 case TGSI_SEMANTIC_TEXCOORD:
115 assert(index < 8);
116 return SI_MAX_IO_GENERIC + 6 + index;
117
118 /* These are rarely used between LS and HS or ES and GS. */
119 case TGSI_SEMANTIC_CLIPDIST:
120 assert(index < 2);
121 return SI_MAX_IO_GENERIC + 6 + 8 + index;
122 case TGSI_SEMANTIC_CLIPVERTEX:
123 return SI_MAX_IO_GENERIC + 6 + 8 + 2;
124 case TGSI_SEMANTIC_PSIZE:
125 return SI_MAX_IO_GENERIC + 6 + 8 + 3;
126
127 /* These can't be written by LS, HS, and ES. */
128 case TGSI_SEMANTIC_LAYER:
129 return SI_MAX_IO_GENERIC + 6 + 8 + 4;
130 case TGSI_SEMANTIC_VIEWPORT_INDEX:
131 return SI_MAX_IO_GENERIC + 6 + 8 + 5;
132 case TGSI_SEMANTIC_PRIMID:
133 STATIC_ASSERT(SI_MAX_IO_GENERIC + 6 + 8 + 6 <= 63);
134 return SI_MAX_IO_GENERIC + 6 + 8 + 6;
135 default:
136 fprintf(stderr, "invalid semantic name = %u\n", semantic_name);
137 assert(!"invalid semantic name");
138 return 0;
139 }
140 }
141
142 static void si_dump_streamout(struct pipe_stream_output_info *so)
143 {
144 unsigned i;
145
146 if (so->num_outputs)
147 fprintf(stderr, "STREAMOUT\n");
148
149 for (i = 0; i < so->num_outputs; i++) {
150 unsigned mask = ((1 << so->output[i].num_components) - 1) << so->output[i].start_component;
151 fprintf(stderr, " %i: BUF%i[%i..%i] <- OUT[%i].%s%s%s%s\n", i, so->output[i].output_buffer,
152 so->output[i].dst_offset, so->output[i].dst_offset + so->output[i].num_components - 1,
153 so->output[i].register_index, mask & 1 ? "x" : "", mask & 2 ? "y" : "",
154 mask & 4 ? "z" : "", mask & 8 ? "w" : "");
155 }
156 }
157
158 static void declare_streamout_params(struct si_shader_context *ctx,
159 struct pipe_stream_output_info *so)
160 {
161 if (ctx->screen->use_ngg_streamout) {
162 if (ctx->type == PIPE_SHADER_TESS_EVAL)
163 ac_add_arg(&ctx->args, AC_ARG_SGPR, 1, AC_ARG_INT, NULL);
164 return;
165 }
166
167 /* Streamout SGPRs. */
168 if (so->num_outputs) {
169 ac_add_arg(&ctx->args, AC_ARG_SGPR, 1, AC_ARG_INT, &ctx->streamout_config);
170 ac_add_arg(&ctx->args, AC_ARG_SGPR, 1, AC_ARG_INT, &ctx->streamout_write_index);
171 } else if (ctx->type == PIPE_SHADER_TESS_EVAL) {
172 ac_add_arg(&ctx->args, AC_ARG_SGPR, 1, AC_ARG_INT, NULL);
173 }
174
175 /* A streamout buffer offset is loaded if the stride is non-zero. */
176 for (int i = 0; i < 4; i++) {
177 if (!so->stride[i])
178 continue;
179
180 ac_add_arg(&ctx->args, AC_ARG_SGPR, 1, AC_ARG_INT, &ctx->streamout_offset[i]);
181 }
182 }
183
184 unsigned si_get_max_workgroup_size(const struct si_shader *shader)
185 {
186 switch (shader->selector->type) {
187 case PIPE_SHADER_VERTEX:
188 case PIPE_SHADER_TESS_EVAL:
189 return shader->key.as_ngg ? 128 : 0;
190
191 case PIPE_SHADER_TESS_CTRL:
192 /* Return this so that LLVM doesn't remove s_barrier
193 * instructions on chips where we use s_barrier. */
194 return shader->selector->screen->info.chip_class >= GFX7 ? 128 : 0;
195
196 case PIPE_SHADER_GEOMETRY:
197 return shader->selector->screen->info.chip_class >= GFX9 ? 128 : 0;
198
199 case PIPE_SHADER_COMPUTE:
200 break; /* see below */
201
202 default:
203 return 0;
204 }
205
206 const unsigned *properties = shader->selector->info.properties;
207 unsigned max_work_group_size = properties[TGSI_PROPERTY_CS_FIXED_BLOCK_WIDTH] *
208 properties[TGSI_PROPERTY_CS_FIXED_BLOCK_HEIGHT] *
209 properties[TGSI_PROPERTY_CS_FIXED_BLOCK_DEPTH];
210
211 if (!max_work_group_size) {
212 /* This is a variable group size compute shader,
213 * compile it for the maximum possible group size.
214 */
215 max_work_group_size = SI_MAX_VARIABLE_THREADS_PER_BLOCK;
216 }
217 return max_work_group_size;
218 }
219
220 static void declare_const_and_shader_buffers(struct si_shader_context *ctx, bool assign_params)
221 {
222 enum ac_arg_type const_shader_buf_type;
223
224 if (ctx->shader->selector->info.const_buffers_declared == 1 &&
225 ctx->shader->selector->info.shader_buffers_declared == 0)
226 const_shader_buf_type = AC_ARG_CONST_FLOAT_PTR;
227 else
228 const_shader_buf_type = AC_ARG_CONST_DESC_PTR;
229
230 ac_add_arg(
231 &ctx->args, AC_ARG_SGPR, 1, const_shader_buf_type,
232 assign_params ? &ctx->const_and_shader_buffers : &ctx->other_const_and_shader_buffers);
233 }
234
235 static void declare_samplers_and_images(struct si_shader_context *ctx, bool assign_params)
236 {
237 ac_add_arg(&ctx->args, AC_ARG_SGPR, 1, AC_ARG_CONST_IMAGE_PTR,
238 assign_params ? &ctx->samplers_and_images : &ctx->other_samplers_and_images);
239 }
240
241 static void declare_per_stage_desc_pointers(struct si_shader_context *ctx, bool assign_params)
242 {
243 declare_const_and_shader_buffers(ctx, assign_params);
244 declare_samplers_and_images(ctx, assign_params);
245 }
246
247 static void declare_global_desc_pointers(struct si_shader_context *ctx)
248 {
249 ac_add_arg(&ctx->args, AC_ARG_SGPR, 1, AC_ARG_CONST_DESC_PTR, &ctx->rw_buffers);
250 ac_add_arg(&ctx->args, AC_ARG_SGPR, 1, AC_ARG_CONST_IMAGE_PTR,
251 &ctx->bindless_samplers_and_images);
252 }
253
254 static void declare_vs_specific_input_sgprs(struct si_shader_context *ctx)
255 {
256 ac_add_arg(&ctx->args, AC_ARG_SGPR, 1, AC_ARG_INT, &ctx->vs_state_bits);
257 if (!ctx->shader->is_gs_copy_shader) {
258 ac_add_arg(&ctx->args, AC_ARG_SGPR, 1, AC_ARG_INT, &ctx->args.base_vertex);
259 ac_add_arg(&ctx->args, AC_ARG_SGPR, 1, AC_ARG_INT, &ctx->args.start_instance);
260 ac_add_arg(&ctx->args, AC_ARG_SGPR, 1, AC_ARG_INT, &ctx->args.draw_id);
261 }
262 }
263
264 static void declare_vb_descriptor_input_sgprs(struct si_shader_context *ctx)
265 {
266 ac_add_arg(&ctx->args, AC_ARG_SGPR, 1, AC_ARG_CONST_DESC_PTR, &ctx->vertex_buffers);
267
268 unsigned num_vbos_in_user_sgprs = ctx->shader->selector->num_vbos_in_user_sgprs;
269 if (num_vbos_in_user_sgprs) {
270 unsigned user_sgprs = ctx->args.num_sgprs_used;
271
272 if (si_is_merged_shader(ctx->shader))
273 user_sgprs -= 8;
274 assert(user_sgprs <= SI_SGPR_VS_VB_DESCRIPTOR_FIRST);
275
276 /* Declare unused SGPRs to align VB descriptors to 4 SGPRs (hw requirement). */
277 for (unsigned i = user_sgprs; i < SI_SGPR_VS_VB_DESCRIPTOR_FIRST; i++)
278 ac_add_arg(&ctx->args, AC_ARG_SGPR, 1, AC_ARG_INT, NULL); /* unused */
279
280 assert(num_vbos_in_user_sgprs <= ARRAY_SIZE(ctx->vb_descriptors));
281 for (unsigned i = 0; i < num_vbos_in_user_sgprs; i++)
282 ac_add_arg(&ctx->args, AC_ARG_SGPR, 4, AC_ARG_INT, &ctx->vb_descriptors[i]);
283 }
284 }
285
286 static void declare_vs_input_vgprs(struct si_shader_context *ctx, unsigned *num_prolog_vgprs,
287 bool ngg_cull_shader)
288 {
289 struct si_shader *shader = ctx->shader;
290
291 ac_add_arg(&ctx->args, AC_ARG_VGPR, 1, AC_ARG_INT, &ctx->args.vertex_id);
292 if (shader->key.as_ls) {
293 ac_add_arg(&ctx->args, AC_ARG_VGPR, 1, AC_ARG_INT, &ctx->rel_auto_id);
294 if (ctx->screen->info.chip_class >= GFX10) {
295 ac_add_arg(&ctx->args, AC_ARG_VGPR, 1, AC_ARG_INT, NULL); /* user VGPR */
296 ac_add_arg(&ctx->args, AC_ARG_VGPR, 1, AC_ARG_INT, &ctx->args.instance_id);
297 } else {
298 ac_add_arg(&ctx->args, AC_ARG_VGPR, 1, AC_ARG_INT, &ctx->args.instance_id);
299 ac_add_arg(&ctx->args, AC_ARG_VGPR, 1, AC_ARG_INT, NULL); /* unused */
300 }
301 } else if (ctx->screen->info.chip_class >= GFX10) {
302 ac_add_arg(&ctx->args, AC_ARG_VGPR, 1, AC_ARG_INT, NULL); /* user VGPR */
303 ac_add_arg(&ctx->args, AC_ARG_VGPR, 1, AC_ARG_INT,
304 &ctx->vs_prim_id); /* user vgpr or PrimID (legacy) */
305 ac_add_arg(&ctx->args, AC_ARG_VGPR, 1, AC_ARG_INT, &ctx->args.instance_id);
306 } else {
307 ac_add_arg(&ctx->args, AC_ARG_VGPR, 1, AC_ARG_INT, &ctx->args.instance_id);
308 ac_add_arg(&ctx->args, AC_ARG_VGPR, 1, AC_ARG_INT, &ctx->vs_prim_id);
309 ac_add_arg(&ctx->args, AC_ARG_VGPR, 1, AC_ARG_INT, NULL); /* unused */
310 }
311
312 if (!shader->is_gs_copy_shader) {
313 if (shader->key.opt.ngg_culling && !ngg_cull_shader) {
314 ac_add_arg(&ctx->args, AC_ARG_VGPR, 1, AC_ARG_INT, &ctx->ngg_old_thread_id);
315 }
316
317 /* Vertex load indices. */
318 if (shader->selector->info.num_inputs) {
319 ac_add_arg(&ctx->args, AC_ARG_VGPR, 1, AC_ARG_INT, &ctx->vertex_index0);
320 for (unsigned i = 1; i < shader->selector->info.num_inputs; i++)
321 ac_add_arg(&ctx->args, AC_ARG_VGPR, 1, AC_ARG_INT, NULL);
322 }
323 *num_prolog_vgprs += shader->selector->info.num_inputs;
324 }
325 }
326
327 static void declare_vs_blit_inputs(struct si_shader_context *ctx, unsigned vs_blit_property)
328 {
329 ac_add_arg(&ctx->args, AC_ARG_SGPR, 1, AC_ARG_INT, &ctx->vs_blit_inputs); /* i16 x1, y1 */
330 ac_add_arg(&ctx->args, AC_ARG_SGPR, 1, AC_ARG_INT, NULL); /* i16 x1, y1 */
331 ac_add_arg(&ctx->args, AC_ARG_SGPR, 1, AC_ARG_FLOAT, NULL); /* depth */
332
333 if (vs_blit_property == SI_VS_BLIT_SGPRS_POS_COLOR) {
334 ac_add_arg(&ctx->args, AC_ARG_SGPR, 1, AC_ARG_FLOAT, NULL); /* color0 */
335 ac_add_arg(&ctx->args, AC_ARG_SGPR, 1, AC_ARG_FLOAT, NULL); /* color1 */
336 ac_add_arg(&ctx->args, AC_ARG_SGPR, 1, AC_ARG_FLOAT, NULL); /* color2 */
337 ac_add_arg(&ctx->args, AC_ARG_SGPR, 1, AC_ARG_FLOAT, NULL); /* color3 */
338 } else if (vs_blit_property == SI_VS_BLIT_SGPRS_POS_TEXCOORD) {
339 ac_add_arg(&ctx->args, AC_ARG_SGPR, 1, AC_ARG_FLOAT, NULL); /* texcoord.x1 */
340 ac_add_arg(&ctx->args, AC_ARG_SGPR, 1, AC_ARG_FLOAT, NULL); /* texcoord.y1 */
341 ac_add_arg(&ctx->args, AC_ARG_SGPR, 1, AC_ARG_FLOAT, NULL); /* texcoord.x2 */
342 ac_add_arg(&ctx->args, AC_ARG_SGPR, 1, AC_ARG_FLOAT, NULL); /* texcoord.y2 */
343 ac_add_arg(&ctx->args, AC_ARG_SGPR, 1, AC_ARG_FLOAT, NULL); /* texcoord.z */
344 ac_add_arg(&ctx->args, AC_ARG_SGPR, 1, AC_ARG_FLOAT, NULL); /* texcoord.w */
345 }
346 }
347
348 static void declare_tes_input_vgprs(struct si_shader_context *ctx, bool ngg_cull_shader)
349 {
350 ac_add_arg(&ctx->args, AC_ARG_VGPR, 1, AC_ARG_FLOAT, &ctx->tes_u);
351 ac_add_arg(&ctx->args, AC_ARG_VGPR, 1, AC_ARG_FLOAT, &ctx->tes_v);
352 ac_add_arg(&ctx->args, AC_ARG_VGPR, 1, AC_ARG_INT, &ctx->tes_rel_patch_id);
353 ac_add_arg(&ctx->args, AC_ARG_VGPR, 1, AC_ARG_INT, &ctx->args.tes_patch_id);
354
355 if (ctx->shader->key.opt.ngg_culling && !ngg_cull_shader) {
356 ac_add_arg(&ctx->args, AC_ARG_VGPR, 1, AC_ARG_INT, &ctx->ngg_old_thread_id);
357 }
358 }
359
360 enum
361 {
362 /* Convenient merged shader definitions. */
363 SI_SHADER_MERGED_VERTEX_TESSCTRL = PIPE_SHADER_TYPES,
364 SI_SHADER_MERGED_VERTEX_OR_TESSEVAL_GEOMETRY,
365 };
366
367 void si_add_arg_checked(struct ac_shader_args *args, enum ac_arg_regfile file, unsigned registers,
368 enum ac_arg_type type, struct ac_arg *arg, unsigned idx)
369 {
370 assert(args->arg_count == idx);
371 ac_add_arg(args, file, registers, type, arg);
372 }
373
374 void si_create_function(struct si_shader_context *ctx, bool ngg_cull_shader)
375 {
376 struct si_shader *shader = ctx->shader;
377 LLVMTypeRef returns[AC_MAX_ARGS];
378 unsigned i, num_return_sgprs;
379 unsigned num_returns = 0;
380 unsigned num_prolog_vgprs = 0;
381 unsigned type = ctx->type;
382 unsigned vs_blit_property = shader->selector->info.properties[TGSI_PROPERTY_VS_BLIT_SGPRS_AMD];
383
384 memset(&ctx->args, 0, sizeof(ctx->args));
385
386 /* Set MERGED shaders. */
387 if (ctx->screen->info.chip_class >= GFX9) {
388 if (shader->key.as_ls || type == PIPE_SHADER_TESS_CTRL)
389 type = SI_SHADER_MERGED_VERTEX_TESSCTRL; /* LS or HS */
390 else if (shader->key.as_es || shader->key.as_ngg || type == PIPE_SHADER_GEOMETRY)
391 type = SI_SHADER_MERGED_VERTEX_OR_TESSEVAL_GEOMETRY;
392 }
393
394 switch (type) {
395 case PIPE_SHADER_VERTEX:
396 declare_global_desc_pointers(ctx);
397
398 if (vs_blit_property) {
399 declare_vs_blit_inputs(ctx, vs_blit_property);
400
401 /* VGPRs */
402 declare_vs_input_vgprs(ctx, &num_prolog_vgprs, ngg_cull_shader);
403 break;
404 }
405
406 declare_per_stage_desc_pointers(ctx, true);
407 declare_vs_specific_input_sgprs(ctx);
408 if (!shader->is_gs_copy_shader)
409 declare_vb_descriptor_input_sgprs(ctx);
410
411 if (shader->key.as_es) {
412 ac_add_arg(&ctx->args, AC_ARG_SGPR, 1, AC_ARG_INT, &ctx->es2gs_offset);
413 } else if (shader->key.as_ls) {
414 /* no extra parameters */
415 } else {
416 /* The locations of the other parameters are assigned dynamically. */
417 declare_streamout_params(ctx, &shader->selector->so);
418 }
419
420 /* VGPRs */
421 declare_vs_input_vgprs(ctx, &num_prolog_vgprs, ngg_cull_shader);
422
423 /* Return values */
424 if (shader->key.opt.vs_as_prim_discard_cs) {
425 for (i = 0; i < 4; i++)
426 returns[num_returns++] = ctx->ac.f32; /* VGPRs */
427 }
428 break;
429
430 case PIPE_SHADER_TESS_CTRL: /* GFX6-GFX8 */
431 declare_global_desc_pointers(ctx);
432 declare_per_stage_desc_pointers(ctx, true);
433 ac_add_arg(&ctx->args, AC_ARG_SGPR, 1, AC_ARG_INT, &ctx->tcs_offchip_layout);
434 ac_add_arg(&ctx->args, AC_ARG_SGPR, 1, AC_ARG_INT, &ctx->tcs_out_lds_offsets);
435 ac_add_arg(&ctx->args, AC_ARG_SGPR, 1, AC_ARG_INT, &ctx->tcs_out_lds_layout);
436 ac_add_arg(&ctx->args, AC_ARG_SGPR, 1, AC_ARG_INT, &ctx->vs_state_bits);
437 ac_add_arg(&ctx->args, AC_ARG_SGPR, 1, AC_ARG_INT, &ctx->tcs_offchip_offset);
438 ac_add_arg(&ctx->args, AC_ARG_SGPR, 1, AC_ARG_INT, &ctx->tcs_factor_offset);
439
440 /* VGPRs */
441 ac_add_arg(&ctx->args, AC_ARG_VGPR, 1, AC_ARG_INT, &ctx->args.tcs_patch_id);
442 ac_add_arg(&ctx->args, AC_ARG_VGPR, 1, AC_ARG_INT, &ctx->args.tcs_rel_ids);
443
444 /* param_tcs_offchip_offset and param_tcs_factor_offset are
445 * placed after the user SGPRs.
446 */
447 for (i = 0; i < GFX6_TCS_NUM_USER_SGPR + 2; i++)
448 returns[num_returns++] = ctx->ac.i32; /* SGPRs */
449 for (i = 0; i < 11; i++)
450 returns[num_returns++] = ctx->ac.f32; /* VGPRs */
451 break;
452
453 case SI_SHADER_MERGED_VERTEX_TESSCTRL:
454 /* Merged stages have 8 system SGPRs at the beginning. */
455 /* SPI_SHADER_USER_DATA_ADDR_LO/HI_HS */
456 declare_per_stage_desc_pointers(ctx, ctx->type == PIPE_SHADER_TESS_CTRL);
457 ac_add_arg(&ctx->args, AC_ARG_SGPR, 1, AC_ARG_INT, &ctx->tcs_offchip_offset);
458 ac_add_arg(&ctx->args, AC_ARG_SGPR, 1, AC_ARG_INT, &ctx->merged_wave_info);
459 ac_add_arg(&ctx->args, AC_ARG_SGPR, 1, AC_ARG_INT, &ctx->tcs_factor_offset);
460 ac_add_arg(&ctx->args, AC_ARG_SGPR, 1, AC_ARG_INT, &ctx->merged_scratch_offset);
461 ac_add_arg(&ctx->args, AC_ARG_SGPR, 1, AC_ARG_INT, NULL); /* unused */
462 ac_add_arg(&ctx->args, AC_ARG_SGPR, 1, AC_ARG_INT, NULL); /* unused */
463
464 declare_global_desc_pointers(ctx);
465 declare_per_stage_desc_pointers(ctx, ctx->type == PIPE_SHADER_VERTEX);
466 declare_vs_specific_input_sgprs(ctx);
467
468 ac_add_arg(&ctx->args, AC_ARG_SGPR, 1, AC_ARG_INT, &ctx->tcs_offchip_layout);
469 ac_add_arg(&ctx->args, AC_ARG_SGPR, 1, AC_ARG_INT, &ctx->tcs_out_lds_offsets);
470 ac_add_arg(&ctx->args, AC_ARG_SGPR, 1, AC_ARG_INT, &ctx->tcs_out_lds_layout);
471 declare_vb_descriptor_input_sgprs(ctx);
472
473 /* VGPRs (first TCS, then VS) */
474 ac_add_arg(&ctx->args, AC_ARG_VGPR, 1, AC_ARG_INT, &ctx->args.tcs_patch_id);
475 ac_add_arg(&ctx->args, AC_ARG_VGPR, 1, AC_ARG_INT, &ctx->args.tcs_rel_ids);
476
477 if (ctx->type == PIPE_SHADER_VERTEX) {
478 declare_vs_input_vgprs(ctx, &num_prolog_vgprs, ngg_cull_shader);
479
480 /* LS return values are inputs to the TCS main shader part. */
481 for (i = 0; i < 8 + GFX9_TCS_NUM_USER_SGPR; i++)
482 returns[num_returns++] = ctx->ac.i32; /* SGPRs */
483 for (i = 0; i < 2; i++)
484 returns[num_returns++] = ctx->ac.f32; /* VGPRs */
485 } else {
486 /* TCS return values are inputs to the TCS epilog.
487 *
488 * param_tcs_offchip_offset, param_tcs_factor_offset,
489 * param_tcs_offchip_layout, and param_rw_buffers
490 * should be passed to the epilog.
491 */
492 for (i = 0; i <= 8 + GFX9_SGPR_TCS_OUT_LAYOUT; i++)
493 returns[num_returns++] = ctx->ac.i32; /* SGPRs */
494 for (i = 0; i < 11; i++)
495 returns[num_returns++] = ctx->ac.f32; /* VGPRs */
496 }
497 break;
498
499 case SI_SHADER_MERGED_VERTEX_OR_TESSEVAL_GEOMETRY:
500 /* Merged stages have 8 system SGPRs at the beginning. */
501 /* SPI_SHADER_USER_DATA_ADDR_LO/HI_GS */
502 declare_per_stage_desc_pointers(ctx, ctx->type == PIPE_SHADER_GEOMETRY);
503
504 if (ctx->shader->key.as_ngg)
505 ac_add_arg(&ctx->args, AC_ARG_SGPR, 1, AC_ARG_INT, &ctx->gs_tg_info);
506 else
507 ac_add_arg(&ctx->args, AC_ARG_SGPR, 1, AC_ARG_INT, &ctx->gs2vs_offset);
508
509 ac_add_arg(&ctx->args, AC_ARG_SGPR, 1, AC_ARG_INT, &ctx->merged_wave_info);
510 ac_add_arg(&ctx->args, AC_ARG_SGPR, 1, AC_ARG_INT, &ctx->tcs_offchip_offset);
511 ac_add_arg(&ctx->args, AC_ARG_SGPR, 1, AC_ARG_INT, &ctx->merged_scratch_offset);
512 ac_add_arg(&ctx->args, AC_ARG_SGPR, 1, AC_ARG_CONST_DESC_PTR,
513 &ctx->small_prim_cull_info); /* SPI_SHADER_PGM_LO_GS << 8 */
514 ac_add_arg(&ctx->args, AC_ARG_SGPR, 1, AC_ARG_INT,
515 NULL); /* unused (SPI_SHADER_PGM_LO/HI_GS >> 24) */
516
517 declare_global_desc_pointers(ctx);
518 if (ctx->type != PIPE_SHADER_VERTEX || !vs_blit_property) {
519 declare_per_stage_desc_pointers(
520 ctx, (ctx->type == PIPE_SHADER_VERTEX || ctx->type == PIPE_SHADER_TESS_EVAL));
521 }
522
523 if (ctx->type == PIPE_SHADER_VERTEX) {
524 if (vs_blit_property)
525 declare_vs_blit_inputs(ctx, vs_blit_property);
526 else
527 declare_vs_specific_input_sgprs(ctx);
528 } else {
529 ac_add_arg(&ctx->args, AC_ARG_SGPR, 1, AC_ARG_INT, &ctx->vs_state_bits);
530 ac_add_arg(&ctx->args, AC_ARG_SGPR, 1, AC_ARG_INT, &ctx->tcs_offchip_layout);
531 ac_add_arg(&ctx->args, AC_ARG_SGPR, 1, AC_ARG_INT, &ctx->tes_offchip_addr);
532 /* Declare as many input SGPRs as the VS has. */
533 }
534
535 if (ctx->type == PIPE_SHADER_VERTEX)
536 declare_vb_descriptor_input_sgprs(ctx);
537
538 /* VGPRs (first GS, then VS/TES) */
539 ac_add_arg(&ctx->args, AC_ARG_VGPR, 1, AC_ARG_INT, &ctx->gs_vtx01_offset);
540 ac_add_arg(&ctx->args, AC_ARG_VGPR, 1, AC_ARG_INT, &ctx->gs_vtx23_offset);
541 ac_add_arg(&ctx->args, AC_ARG_VGPR, 1, AC_ARG_INT, &ctx->args.gs_prim_id);
542 ac_add_arg(&ctx->args, AC_ARG_VGPR, 1, AC_ARG_INT, &ctx->args.gs_invocation_id);
543 ac_add_arg(&ctx->args, AC_ARG_VGPR, 1, AC_ARG_INT, &ctx->gs_vtx45_offset);
544
545 if (ctx->type == PIPE_SHADER_VERTEX) {
546 declare_vs_input_vgprs(ctx, &num_prolog_vgprs, ngg_cull_shader);
547 } else if (ctx->type == PIPE_SHADER_TESS_EVAL) {
548 declare_tes_input_vgprs(ctx, ngg_cull_shader);
549 }
550
551 if ((ctx->shader->key.as_es || ngg_cull_shader) &&
552 (ctx->type == PIPE_SHADER_VERTEX || ctx->type == PIPE_SHADER_TESS_EVAL)) {
553 unsigned num_user_sgprs, num_vgprs;
554
555 if (ctx->type == PIPE_SHADER_VERTEX) {
556 /* For the NGG cull shader, add 1 SGPR to hold
557 * the vertex buffer pointer.
558 */
559 num_user_sgprs = GFX9_VSGS_NUM_USER_SGPR + ngg_cull_shader;
560
561 if (ngg_cull_shader && shader->selector->num_vbos_in_user_sgprs) {
562 assert(num_user_sgprs <= 8 + SI_SGPR_VS_VB_DESCRIPTOR_FIRST);
563 num_user_sgprs =
564 SI_SGPR_VS_VB_DESCRIPTOR_FIRST + shader->selector->num_vbos_in_user_sgprs * 4;
565 }
566 } else {
567 num_user_sgprs = GFX9_TESGS_NUM_USER_SGPR;
568 }
569
570 /* The NGG cull shader has to return all 9 VGPRs + the old thread ID.
571 *
572 * The normal merged ESGS shader only has to return the 5 VGPRs
573 * for the GS stage.
574 */
575 num_vgprs = ngg_cull_shader ? 10 : 5;
576
577 /* ES return values are inputs to GS. */
578 for (i = 0; i < 8 + num_user_sgprs; i++)
579 returns[num_returns++] = ctx->ac.i32; /* SGPRs */
580 for (i = 0; i < num_vgprs; i++)
581 returns[num_returns++] = ctx->ac.f32; /* VGPRs */
582 }
583 break;
584
585 case PIPE_SHADER_TESS_EVAL:
586 declare_global_desc_pointers(ctx);
587 declare_per_stage_desc_pointers(ctx, true);
588 ac_add_arg(&ctx->args, AC_ARG_SGPR, 1, AC_ARG_INT, &ctx->vs_state_bits);
589 ac_add_arg(&ctx->args, AC_ARG_SGPR, 1, AC_ARG_INT, &ctx->tcs_offchip_layout);
590 ac_add_arg(&ctx->args, AC_ARG_SGPR, 1, AC_ARG_INT, &ctx->tes_offchip_addr);
591
592 if (shader->key.as_es) {
593 ac_add_arg(&ctx->args, AC_ARG_SGPR, 1, AC_ARG_INT, &ctx->tcs_offchip_offset);
594 ac_add_arg(&ctx->args, AC_ARG_SGPR, 1, AC_ARG_INT, NULL);
595 ac_add_arg(&ctx->args, AC_ARG_SGPR, 1, AC_ARG_INT, &ctx->es2gs_offset);
596 } else {
597 declare_streamout_params(ctx, &shader->selector->so);
598 ac_add_arg(&ctx->args, AC_ARG_SGPR, 1, AC_ARG_INT, &ctx->tcs_offchip_offset);
599 }
600
601 /* VGPRs */
602 declare_tes_input_vgprs(ctx, ngg_cull_shader);
603 break;
604
605 case PIPE_SHADER_GEOMETRY:
606 declare_global_desc_pointers(ctx);
607 declare_per_stage_desc_pointers(ctx, true);
608 ac_add_arg(&ctx->args, AC_ARG_SGPR, 1, AC_ARG_INT, &ctx->gs2vs_offset);
609 ac_add_arg(&ctx->args, AC_ARG_SGPR, 1, AC_ARG_INT, &ctx->gs_wave_id);
610
611 /* VGPRs */
612 ac_add_arg(&ctx->args, AC_ARG_VGPR, 1, AC_ARG_INT, &ctx->gs_vtx_offset[0]);
613 ac_add_arg(&ctx->args, AC_ARG_VGPR, 1, AC_ARG_INT, &ctx->gs_vtx_offset[1]);
614 ac_add_arg(&ctx->args, AC_ARG_VGPR, 1, AC_ARG_INT, &ctx->args.gs_prim_id);
615 ac_add_arg(&ctx->args, AC_ARG_VGPR, 1, AC_ARG_INT, &ctx->gs_vtx_offset[2]);
616 ac_add_arg(&ctx->args, AC_ARG_VGPR, 1, AC_ARG_INT, &ctx->gs_vtx_offset[3]);
617 ac_add_arg(&ctx->args, AC_ARG_VGPR, 1, AC_ARG_INT, &ctx->gs_vtx_offset[4]);
618 ac_add_arg(&ctx->args, AC_ARG_VGPR, 1, AC_ARG_INT, &ctx->gs_vtx_offset[5]);
619 ac_add_arg(&ctx->args, AC_ARG_VGPR, 1, AC_ARG_INT, &ctx->args.gs_invocation_id);
620 break;
621
622 case PIPE_SHADER_FRAGMENT:
623 declare_global_desc_pointers(ctx);
624 declare_per_stage_desc_pointers(ctx, true);
625 si_add_arg_checked(&ctx->args, AC_ARG_SGPR, 1, AC_ARG_INT, NULL, SI_PARAM_ALPHA_REF);
626 si_add_arg_checked(&ctx->args, AC_ARG_SGPR, 1, AC_ARG_INT, &ctx->args.prim_mask,
627 SI_PARAM_PRIM_MASK);
628
629 si_add_arg_checked(&ctx->args, AC_ARG_VGPR, 2, AC_ARG_INT, &ctx->args.persp_sample,
630 SI_PARAM_PERSP_SAMPLE);
631 si_add_arg_checked(&ctx->args, AC_ARG_VGPR, 2, AC_ARG_INT, &ctx->args.persp_center,
632 SI_PARAM_PERSP_CENTER);
633 si_add_arg_checked(&ctx->args, AC_ARG_VGPR, 2, AC_ARG_INT, &ctx->args.persp_centroid,
634 SI_PARAM_PERSP_CENTROID);
635 si_add_arg_checked(&ctx->args, AC_ARG_VGPR, 3, AC_ARG_INT, NULL, SI_PARAM_PERSP_PULL_MODEL);
636 si_add_arg_checked(&ctx->args, AC_ARG_VGPR, 2, AC_ARG_INT, &ctx->args.linear_sample,
637 SI_PARAM_LINEAR_SAMPLE);
638 si_add_arg_checked(&ctx->args, AC_ARG_VGPR, 2, AC_ARG_INT, &ctx->args.linear_center,
639 SI_PARAM_LINEAR_CENTER);
640 si_add_arg_checked(&ctx->args, AC_ARG_VGPR, 2, AC_ARG_INT, &ctx->args.linear_centroid,
641 SI_PARAM_LINEAR_CENTROID);
642 si_add_arg_checked(&ctx->args, AC_ARG_VGPR, 3, AC_ARG_FLOAT, NULL, SI_PARAM_LINE_STIPPLE_TEX);
643 si_add_arg_checked(&ctx->args, AC_ARG_VGPR, 1, AC_ARG_FLOAT, &ctx->args.frag_pos[0],
644 SI_PARAM_POS_X_FLOAT);
645 si_add_arg_checked(&ctx->args, AC_ARG_VGPR, 1, AC_ARG_FLOAT, &ctx->args.frag_pos[1],
646 SI_PARAM_POS_Y_FLOAT);
647 si_add_arg_checked(&ctx->args, AC_ARG_VGPR, 1, AC_ARG_FLOAT, &ctx->args.frag_pos[2],
648 SI_PARAM_POS_Z_FLOAT);
649 si_add_arg_checked(&ctx->args, AC_ARG_VGPR, 1, AC_ARG_FLOAT, &ctx->args.frag_pos[3],
650 SI_PARAM_POS_W_FLOAT);
651 shader->info.face_vgpr_index = ctx->args.num_vgprs_used;
652 si_add_arg_checked(&ctx->args, AC_ARG_VGPR, 1, AC_ARG_INT, &ctx->args.front_face,
653 SI_PARAM_FRONT_FACE);
654 shader->info.ancillary_vgpr_index = ctx->args.num_vgprs_used;
655 si_add_arg_checked(&ctx->args, AC_ARG_VGPR, 1, AC_ARG_INT, &ctx->args.ancillary,
656 SI_PARAM_ANCILLARY);
657 si_add_arg_checked(&ctx->args, AC_ARG_VGPR, 1, AC_ARG_FLOAT, &ctx->args.sample_coverage,
658 SI_PARAM_SAMPLE_COVERAGE);
659 si_add_arg_checked(&ctx->args, AC_ARG_VGPR, 1, AC_ARG_INT, &ctx->pos_fixed_pt,
660 SI_PARAM_POS_FIXED_PT);
661
662 /* Color inputs from the prolog. */
663 if (shader->selector->info.colors_read) {
664 unsigned num_color_elements = util_bitcount(shader->selector->info.colors_read);
665
666 for (i = 0; i < num_color_elements; i++)
667 ac_add_arg(&ctx->args, AC_ARG_VGPR, 1, AC_ARG_FLOAT, NULL);
668
669 num_prolog_vgprs += num_color_elements;
670 }
671
672 /* Outputs for the epilog. */
673 num_return_sgprs = SI_SGPR_ALPHA_REF + 1;
674 num_returns = num_return_sgprs + util_bitcount(shader->selector->info.colors_written) * 4 +
675 shader->selector->info.writes_z + shader->selector->info.writes_stencil +
676 shader->selector->info.writes_samplemask + 1 /* SampleMaskIn */;
677
678 num_returns = MAX2(num_returns, num_return_sgprs + PS_EPILOG_SAMPLEMASK_MIN_LOC + 1);
679
680 for (i = 0; i < num_return_sgprs; i++)
681 returns[i] = ctx->ac.i32;
682 for (; i < num_returns; i++)
683 returns[i] = ctx->ac.f32;
684 break;
685
686 case PIPE_SHADER_COMPUTE:
687 declare_global_desc_pointers(ctx);
688 declare_per_stage_desc_pointers(ctx, true);
689 if (shader->selector->info.uses_grid_size)
690 ac_add_arg(&ctx->args, AC_ARG_SGPR, 3, AC_ARG_INT, &ctx->args.num_work_groups);
691 if (shader->selector->info.uses_block_size &&
692 shader->selector->info.properties[TGSI_PROPERTY_CS_FIXED_BLOCK_WIDTH] == 0)
693 ac_add_arg(&ctx->args, AC_ARG_SGPR, 3, AC_ARG_INT, &ctx->block_size);
694
695 unsigned cs_user_data_dwords =
696 shader->selector->info.properties[TGSI_PROPERTY_CS_USER_DATA_COMPONENTS_AMD];
697 if (cs_user_data_dwords) {
698 ac_add_arg(&ctx->args, AC_ARG_SGPR, cs_user_data_dwords, AC_ARG_INT, &ctx->cs_user_data);
699 }
700
701 /* Some descriptors can be in user SGPRs. */
702 /* Shader buffers in user SGPRs. */
703 for (unsigned i = 0; i < shader->selector->cs_num_shaderbufs_in_user_sgprs; i++) {
704 while (ctx->args.num_sgprs_used % 4 != 0)
705 ac_add_arg(&ctx->args, AC_ARG_SGPR, 1, AC_ARG_INT, NULL);
706
707 ac_add_arg(&ctx->args, AC_ARG_SGPR, 4, AC_ARG_INT, &ctx->cs_shaderbuf[i]);
708 }
709 /* Images in user SGPRs. */
710 for (unsigned i = 0; i < shader->selector->cs_num_images_in_user_sgprs; i++) {
711 unsigned num_sgprs = shader->selector->info.image_buffers & (1 << i) ? 4 : 8;
712
713 while (ctx->args.num_sgprs_used % num_sgprs != 0)
714 ac_add_arg(&ctx->args, AC_ARG_SGPR, 1, AC_ARG_INT, NULL);
715
716 ac_add_arg(&ctx->args, AC_ARG_SGPR, num_sgprs, AC_ARG_INT, &ctx->cs_image[i]);
717 }
718
719 /* Hardware SGPRs. */
720 for (i = 0; i < 3; i++) {
721 if (shader->selector->info.uses_block_id[i]) {
722 ac_add_arg(&ctx->args, AC_ARG_SGPR, 1, AC_ARG_INT, &ctx->args.workgroup_ids[i]);
723 }
724 }
725 if (shader->selector->info.uses_subgroup_info)
726 ac_add_arg(&ctx->args, AC_ARG_SGPR, 1, AC_ARG_INT, &ctx->args.tg_size);
727
728 /* Hardware VGPRs. */
729 ac_add_arg(&ctx->args, AC_ARG_VGPR, 3, AC_ARG_INT, &ctx->args.local_invocation_ids);
730 break;
731 default:
732 assert(0 && "unimplemented shader");
733 return;
734 }
735
736 si_llvm_create_func(ctx, ngg_cull_shader ? "ngg_cull_main" : "main", returns, num_returns,
737 si_get_max_workgroup_size(shader));
738
739 /* Reserve register locations for VGPR inputs the PS prolog may need. */
740 if (ctx->type == PIPE_SHADER_FRAGMENT && !ctx->shader->is_monolithic) {
741 ac_llvm_add_target_dep_function_attr(
742 ctx->main_fn, "InitialPSInputAddr",
743 S_0286D0_PERSP_SAMPLE_ENA(1) | S_0286D0_PERSP_CENTER_ENA(1) |
744 S_0286D0_PERSP_CENTROID_ENA(1) | S_0286D0_LINEAR_SAMPLE_ENA(1) |
745 S_0286D0_LINEAR_CENTER_ENA(1) | S_0286D0_LINEAR_CENTROID_ENA(1) |
746 S_0286D0_FRONT_FACE_ENA(1) | S_0286D0_ANCILLARY_ENA(1) | S_0286D0_POS_FIXED_PT_ENA(1));
747 }
748
749 shader->info.num_input_sgprs = ctx->args.num_sgprs_used;
750 shader->info.num_input_vgprs = ctx->args.num_vgprs_used;
751
752 assert(shader->info.num_input_vgprs >= num_prolog_vgprs);
753 shader->info.num_input_vgprs -= num_prolog_vgprs;
754
755 if (shader->key.as_ls || ctx->type == PIPE_SHADER_TESS_CTRL) {
756 if (USE_LDS_SYMBOLS && LLVM_VERSION_MAJOR >= 9) {
757 /* The LSHS size is not known until draw time, so we append it
758 * at the end of whatever LDS use there may be in the rest of
759 * the shader (currently none, unless LLVM decides to do its
760 * own LDS-based lowering).
761 */
762 ctx->ac.lds = LLVMAddGlobalInAddressSpace(ctx->ac.module, LLVMArrayType(ctx->ac.i32, 0),
763 "__lds_end", AC_ADDR_SPACE_LDS);
764 LLVMSetAlignment(ctx->ac.lds, 256);
765 } else {
766 ac_declare_lds_as_pointer(&ctx->ac);
767 }
768 }
769
770 /* Unlike radv, we override these arguments in the prolog, so to the
771 * API shader they appear as normal arguments.
772 */
773 if (ctx->type == PIPE_SHADER_VERTEX) {
774 ctx->abi.vertex_id = ac_get_arg(&ctx->ac, ctx->args.vertex_id);
775 ctx->abi.instance_id = ac_get_arg(&ctx->ac, ctx->args.instance_id);
776 } else if (ctx->type == PIPE_SHADER_FRAGMENT) {
777 ctx->abi.persp_centroid = ac_get_arg(&ctx->ac, ctx->args.persp_centroid);
778 ctx->abi.linear_centroid = ac_get_arg(&ctx->ac, ctx->args.linear_centroid);
779 }
780 }
781
782 /* For the UMR disassembler. */
783 #define DEBUGGER_END_OF_CODE_MARKER 0xbf9f0000 /* invalid instruction */
784 #define DEBUGGER_NUM_MARKERS 5
785
786 static bool si_shader_binary_open(struct si_screen *screen, struct si_shader *shader,
787 struct ac_rtld_binary *rtld)
788 {
789 const struct si_shader_selector *sel = shader->selector;
790 const char *part_elfs[5];
791 size_t part_sizes[5];
792 unsigned num_parts = 0;
793
794 #define add_part(shader_or_part) \
795 if (shader_or_part) { \
796 part_elfs[num_parts] = (shader_or_part)->binary.elf_buffer; \
797 part_sizes[num_parts] = (shader_or_part)->binary.elf_size; \
798 num_parts++; \
799 }
800
801 add_part(shader->prolog);
802 add_part(shader->previous_stage);
803 add_part(shader->prolog2);
804 add_part(shader);
805 add_part(shader->epilog);
806
807 #undef add_part
808
809 struct ac_rtld_symbol lds_symbols[2];
810 unsigned num_lds_symbols = 0;
811
812 if (sel && screen->info.chip_class >= GFX9 && !shader->is_gs_copy_shader &&
813 (sel->type == PIPE_SHADER_GEOMETRY || shader->key.as_ngg)) {
814 /* We add this symbol even on LLVM <= 8 to ensure that
815 * shader->config.lds_size is set correctly below.
816 */
817 struct ac_rtld_symbol *sym = &lds_symbols[num_lds_symbols++];
818 sym->name = "esgs_ring";
819 sym->size = shader->gs_info.esgs_ring_size * 4;
820 sym->align = 64 * 1024;
821 }
822
823 if (shader->key.as_ngg && sel->type == PIPE_SHADER_GEOMETRY) {
824 struct ac_rtld_symbol *sym = &lds_symbols[num_lds_symbols++];
825 sym->name = "ngg_emit";
826 sym->size = shader->ngg.ngg_emit_size * 4;
827 sym->align = 4;
828 }
829
830 bool ok = ac_rtld_open(
831 rtld, (struct ac_rtld_open_info){.info = &screen->info,
832 .options =
833 {
834 .halt_at_entry = screen->options.halt_shaders,
835 },
836 .shader_type = tgsi_processor_to_shader_stage(sel->type),
837 .wave_size = si_get_shader_wave_size(shader),
838 .num_parts = num_parts,
839 .elf_ptrs = part_elfs,
840 .elf_sizes = part_sizes,
841 .num_shared_lds_symbols = num_lds_symbols,
842 .shared_lds_symbols = lds_symbols});
843
844 if (rtld->lds_size > 0) {
845 unsigned alloc_granularity = screen->info.chip_class >= GFX7 ? 512 : 256;
846 shader->config.lds_size = align(rtld->lds_size, alloc_granularity) / alloc_granularity;
847 }
848
849 return ok;
850 }
851
852 static unsigned si_get_shader_binary_size(struct si_screen *screen, struct si_shader *shader)
853 {
854 struct ac_rtld_binary rtld;
855 si_shader_binary_open(screen, shader, &rtld);
856 return rtld.exec_size;
857 }
858
859 static bool si_get_external_symbol(void *data, const char *name, uint64_t *value)
860 {
861 uint64_t *scratch_va = data;
862
863 if (!strcmp(scratch_rsrc_dword0_symbol, name)) {
864 *value = (uint32_t)*scratch_va;
865 return true;
866 }
867 if (!strcmp(scratch_rsrc_dword1_symbol, name)) {
868 /* Enable scratch coalescing. */
869 *value = S_008F04_BASE_ADDRESS_HI(*scratch_va >> 32) | S_008F04_SWIZZLE_ENABLE(1);
870 return true;
871 }
872
873 return false;
874 }
875
876 bool si_shader_binary_upload(struct si_screen *sscreen, struct si_shader *shader,
877 uint64_t scratch_va)
878 {
879 struct ac_rtld_binary binary;
880 if (!si_shader_binary_open(sscreen, shader, &binary))
881 return false;
882
883 si_resource_reference(&shader->bo, NULL);
884 shader->bo = si_aligned_buffer_create(
885 &sscreen->b, sscreen->info.cpdma_prefetch_writes_memory ? 0 : SI_RESOURCE_FLAG_READ_ONLY,
886 PIPE_USAGE_IMMUTABLE, align(binary.rx_size, SI_CPDMA_ALIGNMENT), 256);
887 if (!shader->bo)
888 return false;
889
890 /* Upload. */
891 struct ac_rtld_upload_info u = {};
892 u.binary = &binary;
893 u.get_external_symbol = si_get_external_symbol;
894 u.cb_data = &scratch_va;
895 u.rx_va = shader->bo->gpu_address;
896 u.rx_ptr = sscreen->ws->buffer_map(
897 shader->bo->buf, NULL,
898 PIPE_TRANSFER_READ_WRITE | PIPE_TRANSFER_UNSYNCHRONIZED | RADEON_TRANSFER_TEMPORARY);
899 if (!u.rx_ptr)
900 return false;
901
902 bool ok = ac_rtld_upload(&u);
903
904 sscreen->ws->buffer_unmap(shader->bo->buf);
905 ac_rtld_close(&binary);
906
907 return ok;
908 }
909
910 static void si_shader_dump_disassembly(struct si_screen *screen,
911 const struct si_shader_binary *binary,
912 enum pipe_shader_type shader_type, unsigned wave_size,
913 struct pipe_debug_callback *debug, const char *name,
914 FILE *file)
915 {
916 struct ac_rtld_binary rtld_binary;
917
918 if (!ac_rtld_open(&rtld_binary, (struct ac_rtld_open_info){
919 .info = &screen->info,
920 .shader_type = tgsi_processor_to_shader_stage(shader_type),
921 .wave_size = wave_size,
922 .num_parts = 1,
923 .elf_ptrs = &binary->elf_buffer,
924 .elf_sizes = &binary->elf_size}))
925 return;
926
927 const char *disasm;
928 size_t nbytes;
929
930 if (!ac_rtld_get_section_by_name(&rtld_binary, ".AMDGPU.disasm", &disasm, &nbytes))
931 goto out;
932
933 if (nbytes > INT_MAX)
934 goto out;
935
936 if (debug && debug->debug_message) {
937 /* Very long debug messages are cut off, so send the
938 * disassembly one line at a time. This causes more
939 * overhead, but on the plus side it simplifies
940 * parsing of resulting logs.
941 */
942 pipe_debug_message(debug, SHADER_INFO, "Shader Disassembly Begin");
943
944 uint64_t line = 0;
945 while (line < nbytes) {
946 int count = nbytes - line;
947 const char *nl = memchr(disasm + line, '\n', nbytes - line);
948 if (nl)
949 count = nl - (disasm + line);
950
951 if (count) {
952 pipe_debug_message(debug, SHADER_INFO, "%.*s", count, disasm + line);
953 }
954
955 line += count + 1;
956 }
957
958 pipe_debug_message(debug, SHADER_INFO, "Shader Disassembly End");
959 }
960
961 if (file) {
962 fprintf(file, "Shader %s disassembly:\n", name);
963 fprintf(file, "%*s", (int)nbytes, disasm);
964 }
965
966 out:
967 ac_rtld_close(&rtld_binary);
968 }
969
970 static void si_calculate_max_simd_waves(struct si_shader *shader)
971 {
972 struct si_screen *sscreen = shader->selector->screen;
973 struct ac_shader_config *conf = &shader->config;
974 unsigned num_inputs = shader->selector->info.num_inputs;
975 unsigned lds_increment = sscreen->info.chip_class >= GFX7 ? 512 : 256;
976 unsigned lds_per_wave = 0;
977 unsigned max_simd_waves;
978
979 max_simd_waves = sscreen->info.max_wave64_per_simd;
980
981 /* Compute LDS usage for PS. */
982 switch (shader->selector->type) {
983 case PIPE_SHADER_FRAGMENT:
984 /* The minimum usage per wave is (num_inputs * 48). The maximum
985 * usage is (num_inputs * 48 * 16).
986 * We can get anything in between and it varies between waves.
987 *
988 * The 48 bytes per input for a single primitive is equal to
989 * 4 bytes/component * 4 components/input * 3 points.
990 *
991 * Other stages don't know the size at compile time or don't
992 * allocate LDS per wave, but instead they do it per thread group.
993 */
994 lds_per_wave = conf->lds_size * lds_increment + align(num_inputs * 48, lds_increment);
995 break;
996 case PIPE_SHADER_COMPUTE:
997 if (shader->selector) {
998 unsigned max_workgroup_size = si_get_max_workgroup_size(shader);
999 lds_per_wave = (conf->lds_size * lds_increment) /
1000 DIV_ROUND_UP(max_workgroup_size, sscreen->compute_wave_size);
1001 }
1002 break;
1003 default:;
1004 }
1005
1006 /* Compute the per-SIMD wave counts. */
1007 if (conf->num_sgprs) {
1008 max_simd_waves =
1009 MIN2(max_simd_waves, sscreen->info.num_physical_sgprs_per_simd / conf->num_sgprs);
1010 }
1011
1012 if (conf->num_vgprs) {
1013 /* Always print wave limits as Wave64, so that we can compare
1014 * Wave32 and Wave64 with shader-db fairly. */
1015 unsigned max_vgprs = sscreen->info.num_physical_wave64_vgprs_per_simd;
1016 max_simd_waves = MIN2(max_simd_waves, max_vgprs / conf->num_vgprs);
1017 }
1018
1019 unsigned max_lds_per_simd = sscreen->info.lds_size_per_workgroup / 4;
1020 if (lds_per_wave)
1021 max_simd_waves = MIN2(max_simd_waves, max_lds_per_simd / lds_per_wave);
1022
1023 shader->info.max_simd_waves = max_simd_waves;
1024 }
1025
1026 void si_shader_dump_stats_for_shader_db(struct si_screen *screen, struct si_shader *shader,
1027 struct pipe_debug_callback *debug)
1028 {
1029 const struct ac_shader_config *conf = &shader->config;
1030
1031 if (screen->options.debug_disassembly)
1032 si_shader_dump_disassembly(screen, &shader->binary, shader->selector->type,
1033 si_get_shader_wave_size(shader), debug, "main", NULL);
1034
1035 pipe_debug_message(debug, SHADER_INFO,
1036 "Shader Stats: SGPRS: %d VGPRS: %d Code Size: %d "
1037 "LDS: %d Scratch: %d Max Waves: %d Spilled SGPRs: %d "
1038 "Spilled VGPRs: %d PrivMem VGPRs: %d",
1039 conf->num_sgprs, conf->num_vgprs, si_get_shader_binary_size(screen, shader),
1040 conf->lds_size, conf->scratch_bytes_per_wave, shader->info.max_simd_waves,
1041 conf->spilled_sgprs, conf->spilled_vgprs, shader->info.private_mem_vgprs);
1042 }
1043
1044 static void si_shader_dump_stats(struct si_screen *sscreen, struct si_shader *shader, FILE *file,
1045 bool check_debug_option)
1046 {
1047 const struct ac_shader_config *conf = &shader->config;
1048
1049 if (!check_debug_option || si_can_dump_shader(sscreen, shader->selector->type)) {
1050 if (shader->selector->type == PIPE_SHADER_FRAGMENT) {
1051 fprintf(file,
1052 "*** SHADER CONFIG ***\n"
1053 "SPI_PS_INPUT_ADDR = 0x%04x\n"
1054 "SPI_PS_INPUT_ENA = 0x%04x\n",
1055 conf->spi_ps_input_addr, conf->spi_ps_input_ena);
1056 }
1057
1058 fprintf(file,
1059 "*** SHADER STATS ***\n"
1060 "SGPRS: %d\n"
1061 "VGPRS: %d\n"
1062 "Spilled SGPRs: %d\n"
1063 "Spilled VGPRs: %d\n"
1064 "Private memory VGPRs: %d\n"
1065 "Code Size: %d bytes\n"
1066 "LDS: %d blocks\n"
1067 "Scratch: %d bytes per wave\n"
1068 "Max Waves: %d\n"
1069 "********************\n\n\n",
1070 conf->num_sgprs, conf->num_vgprs, conf->spilled_sgprs, conf->spilled_vgprs,
1071 shader->info.private_mem_vgprs, si_get_shader_binary_size(sscreen, shader),
1072 conf->lds_size, conf->scratch_bytes_per_wave, shader->info.max_simd_waves);
1073 }
1074 }
1075
1076 const char *si_get_shader_name(const struct si_shader *shader)
1077 {
1078 switch (shader->selector->type) {
1079 case PIPE_SHADER_VERTEX:
1080 if (shader->key.as_es)
1081 return "Vertex Shader as ES";
1082 else if (shader->key.as_ls)
1083 return "Vertex Shader as LS";
1084 else if (shader->key.opt.vs_as_prim_discard_cs)
1085 return "Vertex Shader as Primitive Discard CS";
1086 else if (shader->key.as_ngg)
1087 return "Vertex Shader as ESGS";
1088 else
1089 return "Vertex Shader as VS";
1090 case PIPE_SHADER_TESS_CTRL:
1091 return "Tessellation Control Shader";
1092 case PIPE_SHADER_TESS_EVAL:
1093 if (shader->key.as_es)
1094 return "Tessellation Evaluation Shader as ES";
1095 else if (shader->key.as_ngg)
1096 return "Tessellation Evaluation Shader as ESGS";
1097 else
1098 return "Tessellation Evaluation Shader as VS";
1099 case PIPE_SHADER_GEOMETRY:
1100 if (shader->is_gs_copy_shader)
1101 return "GS Copy Shader as VS";
1102 else
1103 return "Geometry Shader";
1104 case PIPE_SHADER_FRAGMENT:
1105 return "Pixel Shader";
1106 case PIPE_SHADER_COMPUTE:
1107 return "Compute Shader";
1108 default:
1109 return "Unknown Shader";
1110 }
1111 }
1112
1113 void si_shader_dump(struct si_screen *sscreen, struct si_shader *shader,
1114 struct pipe_debug_callback *debug, FILE *file, bool check_debug_option)
1115 {
1116 enum pipe_shader_type shader_type = shader->selector->type;
1117
1118 if (!check_debug_option || si_can_dump_shader(sscreen, shader_type))
1119 si_dump_shader_key(shader, file);
1120
1121 if (!check_debug_option && shader->binary.llvm_ir_string) {
1122 if (shader->previous_stage && shader->previous_stage->binary.llvm_ir_string) {
1123 fprintf(file, "\n%s - previous stage - LLVM IR:\n\n", si_get_shader_name(shader));
1124 fprintf(file, "%s\n", shader->previous_stage->binary.llvm_ir_string);
1125 }
1126
1127 fprintf(file, "\n%s - main shader part - LLVM IR:\n\n", si_get_shader_name(shader));
1128 fprintf(file, "%s\n", shader->binary.llvm_ir_string);
1129 }
1130
1131 if (!check_debug_option ||
1132 (si_can_dump_shader(sscreen, shader_type) && !(sscreen->debug_flags & DBG(NO_ASM)))) {
1133 unsigned wave_size = si_get_shader_wave_size(shader);
1134
1135 fprintf(file, "\n%s:\n", si_get_shader_name(shader));
1136
1137 if (shader->prolog)
1138 si_shader_dump_disassembly(sscreen, &shader->prolog->binary, shader_type, wave_size, debug,
1139 "prolog", file);
1140 if (shader->previous_stage)
1141 si_shader_dump_disassembly(sscreen, &shader->previous_stage->binary, shader_type,
1142 wave_size, debug, "previous stage", file);
1143 if (shader->prolog2)
1144 si_shader_dump_disassembly(sscreen, &shader->prolog2->binary, shader_type, wave_size,
1145 debug, "prolog2", file);
1146
1147 si_shader_dump_disassembly(sscreen, &shader->binary, shader_type, wave_size, debug, "main",
1148 file);
1149
1150 if (shader->epilog)
1151 si_shader_dump_disassembly(sscreen, &shader->epilog->binary, shader_type, wave_size, debug,
1152 "epilog", file);
1153 fprintf(file, "\n");
1154 }
1155
1156 si_shader_dump_stats(sscreen, shader, file, check_debug_option);
1157 }
1158
1159 static void si_dump_shader_key_vs(const struct si_shader_key *key,
1160 const struct si_vs_prolog_bits *prolog, const char *prefix,
1161 FILE *f)
1162 {
1163 fprintf(f, " %s.instance_divisor_is_one = %u\n", prefix, prolog->instance_divisor_is_one);
1164 fprintf(f, " %s.instance_divisor_is_fetched = %u\n", prefix,
1165 prolog->instance_divisor_is_fetched);
1166 fprintf(f, " %s.unpack_instance_id_from_vertex_id = %u\n", prefix,
1167 prolog->unpack_instance_id_from_vertex_id);
1168 fprintf(f, " %s.ls_vgpr_fix = %u\n", prefix, prolog->ls_vgpr_fix);
1169
1170 fprintf(f, " mono.vs.fetch_opencode = %x\n", key->mono.vs_fetch_opencode);
1171 fprintf(f, " mono.vs.fix_fetch = {");
1172 for (int i = 0; i < SI_MAX_ATTRIBS; i++) {
1173 union si_vs_fix_fetch fix = key->mono.vs_fix_fetch[i];
1174 if (i)
1175 fprintf(f, ", ");
1176 if (!fix.bits)
1177 fprintf(f, "0");
1178 else
1179 fprintf(f, "%u.%u.%u.%u", fix.u.reverse, fix.u.log_size, fix.u.num_channels_m1,
1180 fix.u.format);
1181 }
1182 fprintf(f, "}\n");
1183 }
1184
1185 static void si_dump_shader_key(const struct si_shader *shader, FILE *f)
1186 {
1187 const struct si_shader_key *key = &shader->key;
1188 enum pipe_shader_type shader_type = shader->selector->type;
1189
1190 fprintf(f, "SHADER KEY\n");
1191
1192 switch (shader_type) {
1193 case PIPE_SHADER_VERTEX:
1194 si_dump_shader_key_vs(key, &key->part.vs.prolog, "part.vs.prolog", f);
1195 fprintf(f, " as_es = %u\n", key->as_es);
1196 fprintf(f, " as_ls = %u\n", key->as_ls);
1197 fprintf(f, " as_ngg = %u\n", key->as_ngg);
1198 fprintf(f, " mono.u.vs_export_prim_id = %u\n", key->mono.u.vs_export_prim_id);
1199 fprintf(f, " opt.vs_as_prim_discard_cs = %u\n", key->opt.vs_as_prim_discard_cs);
1200 fprintf(f, " opt.cs_prim_type = %s\n", tgsi_primitive_names[key->opt.cs_prim_type]);
1201 fprintf(f, " opt.cs_indexed = %u\n", key->opt.cs_indexed);
1202 fprintf(f, " opt.cs_instancing = %u\n", key->opt.cs_instancing);
1203 fprintf(f, " opt.cs_primitive_restart = %u\n", key->opt.cs_primitive_restart);
1204 fprintf(f, " opt.cs_provoking_vertex_first = %u\n", key->opt.cs_provoking_vertex_first);
1205 fprintf(f, " opt.cs_need_correct_orientation = %u\n", key->opt.cs_need_correct_orientation);
1206 fprintf(f, " opt.cs_cull_front = %u\n", key->opt.cs_cull_front);
1207 fprintf(f, " opt.cs_cull_back = %u\n", key->opt.cs_cull_back);
1208 fprintf(f, " opt.cs_cull_z = %u\n", key->opt.cs_cull_z);
1209 fprintf(f, " opt.cs_halfz_clip_space = %u\n", key->opt.cs_halfz_clip_space);
1210 break;
1211
1212 case PIPE_SHADER_TESS_CTRL:
1213 if (shader->selector->screen->info.chip_class >= GFX9) {
1214 si_dump_shader_key_vs(key, &key->part.tcs.ls_prolog, "part.tcs.ls_prolog", f);
1215 }
1216 fprintf(f, " part.tcs.epilog.prim_mode = %u\n", key->part.tcs.epilog.prim_mode);
1217 fprintf(f, " mono.u.ff_tcs_inputs_to_copy = 0x%" PRIx64 "\n",
1218 key->mono.u.ff_tcs_inputs_to_copy);
1219 break;
1220
1221 case PIPE_SHADER_TESS_EVAL:
1222 fprintf(f, " as_es = %u\n", key->as_es);
1223 fprintf(f, " as_ngg = %u\n", key->as_ngg);
1224 fprintf(f, " mono.u.vs_export_prim_id = %u\n", key->mono.u.vs_export_prim_id);
1225 break;
1226
1227 case PIPE_SHADER_GEOMETRY:
1228 if (shader->is_gs_copy_shader)
1229 break;
1230
1231 if (shader->selector->screen->info.chip_class >= GFX9 &&
1232 key->part.gs.es->type == PIPE_SHADER_VERTEX) {
1233 si_dump_shader_key_vs(key, &key->part.gs.vs_prolog, "part.gs.vs_prolog", f);
1234 }
1235 fprintf(f, " part.gs.prolog.tri_strip_adj_fix = %u\n",
1236 key->part.gs.prolog.tri_strip_adj_fix);
1237 fprintf(f, " part.gs.prolog.gfx9_prev_is_vs = %u\n", key->part.gs.prolog.gfx9_prev_is_vs);
1238 fprintf(f, " as_ngg = %u\n", key->as_ngg);
1239 break;
1240
1241 case PIPE_SHADER_COMPUTE:
1242 break;
1243
1244 case PIPE_SHADER_FRAGMENT:
1245 fprintf(f, " part.ps.prolog.color_two_side = %u\n", key->part.ps.prolog.color_two_side);
1246 fprintf(f, " part.ps.prolog.flatshade_colors = %u\n", key->part.ps.prolog.flatshade_colors);
1247 fprintf(f, " part.ps.prolog.poly_stipple = %u\n", key->part.ps.prolog.poly_stipple);
1248 fprintf(f, " part.ps.prolog.force_persp_sample_interp = %u\n",
1249 key->part.ps.prolog.force_persp_sample_interp);
1250 fprintf(f, " part.ps.prolog.force_linear_sample_interp = %u\n",
1251 key->part.ps.prolog.force_linear_sample_interp);
1252 fprintf(f, " part.ps.prolog.force_persp_center_interp = %u\n",
1253 key->part.ps.prolog.force_persp_center_interp);
1254 fprintf(f, " part.ps.prolog.force_linear_center_interp = %u\n",
1255 key->part.ps.prolog.force_linear_center_interp);
1256 fprintf(f, " part.ps.prolog.bc_optimize_for_persp = %u\n",
1257 key->part.ps.prolog.bc_optimize_for_persp);
1258 fprintf(f, " part.ps.prolog.bc_optimize_for_linear = %u\n",
1259 key->part.ps.prolog.bc_optimize_for_linear);
1260 fprintf(f, " part.ps.prolog.samplemask_log_ps_iter = %u\n",
1261 key->part.ps.prolog.samplemask_log_ps_iter);
1262 fprintf(f, " part.ps.epilog.spi_shader_col_format = 0x%x\n",
1263 key->part.ps.epilog.spi_shader_col_format);
1264 fprintf(f, " part.ps.epilog.color_is_int8 = 0x%X\n", key->part.ps.epilog.color_is_int8);
1265 fprintf(f, " part.ps.epilog.color_is_int10 = 0x%X\n", key->part.ps.epilog.color_is_int10);
1266 fprintf(f, " part.ps.epilog.last_cbuf = %u\n", key->part.ps.epilog.last_cbuf);
1267 fprintf(f, " part.ps.epilog.alpha_func = %u\n", key->part.ps.epilog.alpha_func);
1268 fprintf(f, " part.ps.epilog.alpha_to_one = %u\n", key->part.ps.epilog.alpha_to_one);
1269 fprintf(f, " part.ps.epilog.poly_line_smoothing = %u\n",
1270 key->part.ps.epilog.poly_line_smoothing);
1271 fprintf(f, " part.ps.epilog.clamp_color = %u\n", key->part.ps.epilog.clamp_color);
1272 fprintf(f, " mono.u.ps.interpolate_at_sample_force_center = %u\n",
1273 key->mono.u.ps.interpolate_at_sample_force_center);
1274 fprintf(f, " mono.u.ps.fbfetch_msaa = %u\n", key->mono.u.ps.fbfetch_msaa);
1275 fprintf(f, " mono.u.ps.fbfetch_is_1D = %u\n", key->mono.u.ps.fbfetch_is_1D);
1276 fprintf(f, " mono.u.ps.fbfetch_layered = %u\n", key->mono.u.ps.fbfetch_layered);
1277 break;
1278
1279 default:
1280 assert(0);
1281 }
1282
1283 if ((shader_type == PIPE_SHADER_GEOMETRY || shader_type == PIPE_SHADER_TESS_EVAL ||
1284 shader_type == PIPE_SHADER_VERTEX) &&
1285 !key->as_es && !key->as_ls) {
1286 fprintf(f, " opt.kill_outputs = 0x%" PRIx64 "\n", key->opt.kill_outputs);
1287 fprintf(f, " opt.clip_disable = %u\n", key->opt.clip_disable);
1288 if (shader_type != PIPE_SHADER_GEOMETRY)
1289 fprintf(f, " opt.ngg_culling = 0x%x\n", key->opt.ngg_culling);
1290 }
1291 }
1292
1293 static void si_optimize_vs_outputs(struct si_shader_context *ctx)
1294 {
1295 struct si_shader *shader = ctx->shader;
1296 struct si_shader_info *info = &shader->selector->info;
1297 unsigned skip_vs_optim_mask = 0;
1298
1299 if ((ctx->type != PIPE_SHADER_VERTEX && ctx->type != PIPE_SHADER_TESS_EVAL) ||
1300 shader->key.as_ls || shader->key.as_es)
1301 return;
1302
1303 /* Optimizing these outputs is not possible, since they might be overriden
1304 * at runtime with S_028644_PT_SPRITE_TEX. */
1305 for (int i = 0; i < info->num_outputs; i++) {
1306 if (info->output_semantic_name[i] == TGSI_SEMANTIC_PCOORD ||
1307 info->output_semantic_name[i] == TGSI_SEMANTIC_TEXCOORD) {
1308 skip_vs_optim_mask |= 1u << shader->info.vs_output_param_offset[i];
1309 }
1310 }
1311
1312 ac_optimize_vs_outputs(&ctx->ac, ctx->main_fn, shader->info.vs_output_param_offset,
1313 info->num_outputs, skip_vs_optim_mask,
1314 &shader->info.nr_param_exports);
1315 }
1316
1317 static bool si_vs_needs_prolog(const struct si_shader_selector *sel,
1318 const struct si_vs_prolog_bits *prolog_key,
1319 const struct si_shader_key *key, bool ngg_cull_shader)
1320 {
1321 /* VGPR initialization fixup for Vega10 and Raven is always done in the
1322 * VS prolog. */
1323 return sel->vs_needs_prolog || prolog_key->ls_vgpr_fix ||
1324 prolog_key->unpack_instance_id_from_vertex_id ||
1325 (ngg_cull_shader && key->opt.ngg_culling & SI_NGG_CULL_GS_FAST_LAUNCH_ALL);
1326 }
1327
1328 static bool si_build_main_function(struct si_shader_context *ctx, struct si_shader *shader,
1329 struct nir_shader *nir, bool free_nir, bool ngg_cull_shader)
1330 {
1331 struct si_shader_selector *sel = shader->selector;
1332 const struct si_shader_info *info = &sel->info;
1333
1334 ctx->shader = shader;
1335 ctx->type = sel->type;
1336
1337 ctx->num_const_buffers = util_last_bit(info->const_buffers_declared);
1338 ctx->num_shader_buffers = util_last_bit(info->shader_buffers_declared);
1339
1340 ctx->num_samplers = util_last_bit(info->samplers_declared);
1341 ctx->num_images = util_last_bit(info->images_declared);
1342
1343 si_llvm_init_resource_callbacks(ctx);
1344
1345 switch (ctx->type) {
1346 case PIPE_SHADER_VERTEX:
1347 si_llvm_init_vs_callbacks(ctx, ngg_cull_shader);
1348 break;
1349 case PIPE_SHADER_TESS_CTRL:
1350 si_llvm_init_tcs_callbacks(ctx);
1351 break;
1352 case PIPE_SHADER_TESS_EVAL:
1353 si_llvm_init_tes_callbacks(ctx, ngg_cull_shader);
1354 break;
1355 case PIPE_SHADER_GEOMETRY:
1356 si_llvm_init_gs_callbacks(ctx);
1357 break;
1358 case PIPE_SHADER_FRAGMENT:
1359 si_llvm_init_ps_callbacks(ctx);
1360 break;
1361 case PIPE_SHADER_COMPUTE:
1362 ctx->abi.load_local_group_size = si_llvm_get_block_size;
1363 break;
1364 default:
1365 assert(!"Unsupported shader type");
1366 return false;
1367 }
1368
1369 si_create_function(ctx, ngg_cull_shader);
1370
1371 if (ctx->shader->key.as_es || ctx->type == PIPE_SHADER_GEOMETRY)
1372 si_preload_esgs_ring(ctx);
1373
1374 if (ctx->type == PIPE_SHADER_GEOMETRY)
1375 si_preload_gs_rings(ctx);
1376 else if (ctx->type == PIPE_SHADER_TESS_EVAL)
1377 si_llvm_preload_tes_rings(ctx);
1378
1379 if (ctx->type == PIPE_SHADER_TESS_CTRL && sel->info.tessfactors_are_def_in_all_invocs) {
1380 for (unsigned i = 0; i < 6; i++) {
1381 ctx->invoc0_tess_factors[i] = ac_build_alloca_undef(&ctx->ac, ctx->ac.i32, "");
1382 }
1383 }
1384
1385 if (ctx->type == PIPE_SHADER_GEOMETRY) {
1386 for (unsigned i = 0; i < 4; i++) {
1387 ctx->gs_next_vertex[i] = ac_build_alloca(&ctx->ac, ctx->ac.i32, "");
1388 }
1389 if (shader->key.as_ngg) {
1390 for (unsigned i = 0; i < 4; ++i) {
1391 ctx->gs_curprim_verts[i] = ac_build_alloca(&ctx->ac, ctx->ac.i32, "");
1392 ctx->gs_generated_prims[i] = ac_build_alloca(&ctx->ac, ctx->ac.i32, "");
1393 }
1394
1395 assert(!ctx->gs_ngg_scratch);
1396 LLVMTypeRef ai32 = LLVMArrayType(ctx->ac.i32, gfx10_ngg_get_scratch_dw_size(shader));
1397 ctx->gs_ngg_scratch =
1398 LLVMAddGlobalInAddressSpace(ctx->ac.module, ai32, "ngg_scratch", AC_ADDR_SPACE_LDS);
1399 LLVMSetInitializer(ctx->gs_ngg_scratch, LLVMGetUndef(ai32));
1400 LLVMSetAlignment(ctx->gs_ngg_scratch, 4);
1401
1402 ctx->gs_ngg_emit = LLVMAddGlobalInAddressSpace(
1403 ctx->ac.module, LLVMArrayType(ctx->ac.i32, 0), "ngg_emit", AC_ADDR_SPACE_LDS);
1404 LLVMSetLinkage(ctx->gs_ngg_emit, LLVMExternalLinkage);
1405 LLVMSetAlignment(ctx->gs_ngg_emit, 4);
1406 }
1407 }
1408
1409 if (ctx->type != PIPE_SHADER_GEOMETRY && (shader->key.as_ngg && !shader->key.as_es)) {
1410 /* Unconditionally declare scratch space base for streamout and
1411 * vertex compaction. Whether space is actually allocated is
1412 * determined during linking / PM4 creation.
1413 *
1414 * Add an extra dword per vertex to ensure an odd stride, which
1415 * avoids bank conflicts for SoA accesses.
1416 */
1417 if (!gfx10_is_ngg_passthrough(shader))
1418 si_llvm_declare_esgs_ring(ctx);
1419
1420 /* This is really only needed when streamout and / or vertex
1421 * compaction is enabled.
1422 */
1423 if (!ctx->gs_ngg_scratch && (sel->so.num_outputs || shader->key.opt.ngg_culling)) {
1424 LLVMTypeRef asi32 = LLVMArrayType(ctx->ac.i32, gfx10_ngg_get_scratch_dw_size(shader));
1425 ctx->gs_ngg_scratch =
1426 LLVMAddGlobalInAddressSpace(ctx->ac.module, asi32, "ngg_scratch", AC_ADDR_SPACE_LDS);
1427 LLVMSetInitializer(ctx->gs_ngg_scratch, LLVMGetUndef(asi32));
1428 LLVMSetAlignment(ctx->gs_ngg_scratch, 4);
1429 }
1430 }
1431
1432 /* For GFX9 merged shaders:
1433 * - Set EXEC for the first shader. If the prolog is present, set
1434 * EXEC there instead.
1435 * - Add a barrier before the second shader.
1436 * - In the second shader, reset EXEC to ~0 and wrap the main part in
1437 * an if-statement. This is required for correctness in geometry
1438 * shaders, to ensure that empty GS waves do not send GS_EMIT and
1439 * GS_CUT messages.
1440 *
1441 * For monolithic merged shaders, the first shader is wrapped in an
1442 * if-block together with its prolog in si_build_wrapper_function.
1443 *
1444 * NGG vertex and tess eval shaders running as the last
1445 * vertex/geometry stage handle execution explicitly using
1446 * if-statements.
1447 */
1448 if (ctx->screen->info.chip_class >= GFX9) {
1449 if (!shader->is_monolithic && (shader->key.as_es || shader->key.as_ls) &&
1450 (ctx->type == PIPE_SHADER_TESS_EVAL ||
1451 (ctx->type == PIPE_SHADER_VERTEX &&
1452 !si_vs_needs_prolog(sel, &shader->key.part.vs.prolog, &shader->key, ngg_cull_shader)))) {
1453 si_init_exec_from_input(ctx, ctx->merged_wave_info, 0);
1454 } else if (ctx->type == PIPE_SHADER_TESS_CTRL || ctx->type == PIPE_SHADER_GEOMETRY ||
1455 (shader->key.as_ngg && !shader->key.as_es)) {
1456 LLVMValueRef thread_enabled;
1457 bool nested_barrier;
1458
1459 if (!shader->is_monolithic || (ctx->type == PIPE_SHADER_TESS_EVAL && shader->key.as_ngg &&
1460 !shader->key.as_es && !shader->key.opt.ngg_culling))
1461 ac_init_exec_full_mask(&ctx->ac);
1462
1463 if ((ctx->type == PIPE_SHADER_VERTEX || ctx->type == PIPE_SHADER_TESS_EVAL) &&
1464 shader->key.as_ngg && !shader->key.as_es && !shader->key.opt.ngg_culling) {
1465 gfx10_ngg_build_sendmsg_gs_alloc_req(ctx);
1466
1467 /* Build the primitive export at the beginning
1468 * of the shader if possible.
1469 */
1470 if (gfx10_ngg_export_prim_early(shader))
1471 gfx10_ngg_build_export_prim(ctx, NULL, NULL);
1472 }
1473
1474 if (ctx->type == PIPE_SHADER_TESS_CTRL || ctx->type == PIPE_SHADER_GEOMETRY) {
1475 if (ctx->type == PIPE_SHADER_GEOMETRY && shader->key.as_ngg) {
1476 gfx10_ngg_gs_emit_prologue(ctx);
1477 nested_barrier = false;
1478 } else {
1479 nested_barrier = true;
1480 }
1481
1482 thread_enabled = si_is_gs_thread(ctx);
1483 } else {
1484 thread_enabled = si_is_es_thread(ctx);
1485 nested_barrier = false;
1486 }
1487
1488 ctx->merged_wrap_if_entry_block = LLVMGetInsertBlock(ctx->ac.builder);
1489 ctx->merged_wrap_if_label = 11500;
1490 ac_build_ifcc(&ctx->ac, thread_enabled, ctx->merged_wrap_if_label);
1491
1492 if (nested_barrier) {
1493 /* Execute a barrier before the second shader in
1494 * a merged shader.
1495 *
1496 * Execute the barrier inside the conditional block,
1497 * so that empty waves can jump directly to s_endpgm,
1498 * which will also signal the barrier.
1499 *
1500 * This is possible in gfx9, because an empty wave
1501 * for the second shader does not participate in
1502 * the epilogue. With NGG, empty waves may still
1503 * be required to export data (e.g. GS output vertices),
1504 * so we cannot let them exit early.
1505 *
1506 * If the shader is TCS and the TCS epilog is present
1507 * and contains a barrier, it will wait there and then
1508 * reach s_endpgm.
1509 */
1510 si_llvm_emit_barrier(ctx);
1511 }
1512 }
1513 }
1514
1515 bool success = si_nir_build_llvm(ctx, nir);
1516 if (free_nir)
1517 ralloc_free(nir);
1518 if (!success) {
1519 fprintf(stderr, "Failed to translate shader from NIR to LLVM\n");
1520 return false;
1521 }
1522
1523 si_llvm_build_ret(ctx, ctx->return_value);
1524 return true;
1525 }
1526
1527 /**
1528 * Compute the VS prolog key, which contains all the information needed to
1529 * build the VS prolog function, and set shader->info bits where needed.
1530 *
1531 * \param info Shader info of the vertex shader.
1532 * \param num_input_sgprs Number of input SGPRs for the vertex shader.
1533 * \param has_old_ Whether the preceding shader part is the NGG cull shader.
1534 * \param prolog_key Key of the VS prolog
1535 * \param shader_out The vertex shader, or the next shader if merging LS+HS or ES+GS.
1536 * \param key Output shader part key.
1537 */
1538 static void si_get_vs_prolog_key(const struct si_shader_info *info, unsigned num_input_sgprs,
1539 bool ngg_cull_shader, const struct si_vs_prolog_bits *prolog_key,
1540 struct si_shader *shader_out, union si_shader_part_key *key)
1541 {
1542 memset(key, 0, sizeof(*key));
1543 key->vs_prolog.states = *prolog_key;
1544 key->vs_prolog.num_input_sgprs = num_input_sgprs;
1545 key->vs_prolog.num_inputs = info->num_inputs;
1546 key->vs_prolog.as_ls = shader_out->key.as_ls;
1547 key->vs_prolog.as_es = shader_out->key.as_es;
1548 key->vs_prolog.as_ngg = shader_out->key.as_ngg;
1549 key->vs_prolog.as_prim_discard_cs = shader_out->key.opt.vs_as_prim_discard_cs;
1550
1551 if (ngg_cull_shader) {
1552 key->vs_prolog.gs_fast_launch_tri_list =
1553 !!(shader_out->key.opt.ngg_culling & SI_NGG_CULL_GS_FAST_LAUNCH_TRI_LIST);
1554 key->vs_prolog.gs_fast_launch_tri_strip =
1555 !!(shader_out->key.opt.ngg_culling & SI_NGG_CULL_GS_FAST_LAUNCH_TRI_STRIP);
1556 } else {
1557 key->vs_prolog.has_ngg_cull_inputs = !!shader_out->key.opt.ngg_culling;
1558 }
1559
1560 if (shader_out->selector->type == PIPE_SHADER_TESS_CTRL) {
1561 key->vs_prolog.as_ls = 1;
1562 key->vs_prolog.num_merged_next_stage_vgprs = 2;
1563 } else if (shader_out->selector->type == PIPE_SHADER_GEOMETRY) {
1564 key->vs_prolog.as_es = 1;
1565 key->vs_prolog.num_merged_next_stage_vgprs = 5;
1566 } else if (shader_out->key.as_ngg) {
1567 key->vs_prolog.num_merged_next_stage_vgprs = 5;
1568 }
1569
1570 /* Only one of these combinations can be set. as_ngg can be set with as_es. */
1571 assert(key->vs_prolog.as_ls + key->vs_prolog.as_ngg +
1572 (key->vs_prolog.as_es && !key->vs_prolog.as_ngg) + key->vs_prolog.as_prim_discard_cs <=
1573 1);
1574
1575 /* Enable loading the InstanceID VGPR. */
1576 uint16_t input_mask = u_bit_consecutive(0, info->num_inputs);
1577
1578 if ((key->vs_prolog.states.instance_divisor_is_one |
1579 key->vs_prolog.states.instance_divisor_is_fetched) &
1580 input_mask)
1581 shader_out->info.uses_instanceid = true;
1582 }
1583
1584 static bool si_should_optimize_less(struct ac_llvm_compiler *compiler,
1585 struct si_shader_selector *sel)
1586 {
1587 if (!compiler->low_opt_passes)
1588 return false;
1589
1590 /* Assume a slow CPU. */
1591 assert(!sel->screen->info.has_dedicated_vram && sel->screen->info.chip_class <= GFX8);
1592
1593 /* For a crazy dEQP test containing 2597 memory opcodes, mostly
1594 * buffer stores. */
1595 return sel->type == PIPE_SHADER_COMPUTE && sel->info.num_memory_instructions > 1000;
1596 }
1597
1598 static struct nir_shader *get_nir_shader(struct si_shader_selector *sel, bool *free_nir)
1599 {
1600 *free_nir = false;
1601
1602 if (sel->nir) {
1603 return sel->nir;
1604 } else if (sel->nir_binary) {
1605 struct pipe_screen *screen = &sel->screen->b;
1606 const void *options = screen->get_compiler_options(screen, PIPE_SHADER_IR_NIR, sel->type);
1607
1608 struct blob_reader blob_reader;
1609 blob_reader_init(&blob_reader, sel->nir_binary, sel->nir_size);
1610 *free_nir = true;
1611 return nir_deserialize(NULL, options, &blob_reader);
1612 }
1613 return NULL;
1614 }
1615
1616 static bool si_llvm_compile_shader(struct si_screen *sscreen, struct ac_llvm_compiler *compiler,
1617 struct si_shader *shader, struct pipe_debug_callback *debug,
1618 struct nir_shader *nir, bool free_nir)
1619 {
1620 struct si_shader_selector *sel = shader->selector;
1621 struct si_shader_context ctx;
1622
1623 si_llvm_context_init(&ctx, sscreen, compiler, si_get_shader_wave_size(shader));
1624
1625 LLVMValueRef ngg_cull_main_fn = NULL;
1626 if (shader->key.opt.ngg_culling) {
1627 if (!si_build_main_function(&ctx, shader, nir, false, true)) {
1628 si_llvm_dispose(&ctx);
1629 return false;
1630 }
1631 ngg_cull_main_fn = ctx.main_fn;
1632 ctx.main_fn = NULL;
1633 }
1634
1635 if (!si_build_main_function(&ctx, shader, nir, free_nir, false)) {
1636 si_llvm_dispose(&ctx);
1637 return false;
1638 }
1639
1640 if (shader->is_monolithic && ctx.type == PIPE_SHADER_VERTEX) {
1641 LLVMValueRef parts[4];
1642 unsigned num_parts = 0;
1643 bool has_prolog = false;
1644 LLVMValueRef main_fn = ctx.main_fn;
1645
1646 if (ngg_cull_main_fn) {
1647 if (si_vs_needs_prolog(sel, &shader->key.part.vs.prolog, &shader->key, true)) {
1648 union si_shader_part_key prolog_key;
1649 si_get_vs_prolog_key(&sel->info, shader->info.num_input_sgprs, true,
1650 &shader->key.part.vs.prolog, shader, &prolog_key);
1651 prolog_key.vs_prolog.is_monolithic = true;
1652 si_llvm_build_vs_prolog(&ctx, &prolog_key);
1653 parts[num_parts++] = ctx.main_fn;
1654 has_prolog = true;
1655 }
1656 parts[num_parts++] = ngg_cull_main_fn;
1657 }
1658
1659 if (si_vs_needs_prolog(sel, &shader->key.part.vs.prolog, &shader->key, false)) {
1660 union si_shader_part_key prolog_key;
1661 si_get_vs_prolog_key(&sel->info, shader->info.num_input_sgprs, false,
1662 &shader->key.part.vs.prolog, shader, &prolog_key);
1663 prolog_key.vs_prolog.is_monolithic = true;
1664 si_llvm_build_vs_prolog(&ctx, &prolog_key);
1665 parts[num_parts++] = ctx.main_fn;
1666 has_prolog = true;
1667 }
1668 parts[num_parts++] = main_fn;
1669
1670 si_build_wrapper_function(&ctx, parts, num_parts, has_prolog ? 1 : 0, 0);
1671
1672 if (ctx.shader->key.opt.vs_as_prim_discard_cs)
1673 si_build_prim_discard_compute_shader(&ctx);
1674 } else if (shader->is_monolithic && ctx.type == PIPE_SHADER_TESS_EVAL && ngg_cull_main_fn) {
1675 LLVMValueRef parts[2];
1676
1677 parts[0] = ngg_cull_main_fn;
1678 parts[1] = ctx.main_fn;
1679
1680 si_build_wrapper_function(&ctx, parts, 2, 0, 0);
1681 } else if (shader->is_monolithic && ctx.type == PIPE_SHADER_TESS_CTRL) {
1682 if (sscreen->info.chip_class >= GFX9) {
1683 struct si_shader_selector *ls = shader->key.part.tcs.ls;
1684 LLVMValueRef parts[4];
1685 bool vs_needs_prolog =
1686 si_vs_needs_prolog(ls, &shader->key.part.tcs.ls_prolog, &shader->key, false);
1687
1688 /* TCS main part */
1689 parts[2] = ctx.main_fn;
1690
1691 /* TCS epilog */
1692 union si_shader_part_key tcs_epilog_key;
1693 memset(&tcs_epilog_key, 0, sizeof(tcs_epilog_key));
1694 tcs_epilog_key.tcs_epilog.states = shader->key.part.tcs.epilog;
1695 si_llvm_build_tcs_epilog(&ctx, &tcs_epilog_key);
1696 parts[3] = ctx.main_fn;
1697
1698 /* VS as LS main part */
1699 nir = get_nir_shader(ls, &free_nir);
1700 struct si_shader shader_ls = {};
1701 shader_ls.selector = ls;
1702 shader_ls.key.as_ls = 1;
1703 shader_ls.key.mono = shader->key.mono;
1704 shader_ls.key.opt = shader->key.opt;
1705 shader_ls.is_monolithic = true;
1706
1707 if (!si_build_main_function(&ctx, &shader_ls, nir, free_nir, false)) {
1708 si_llvm_dispose(&ctx);
1709 return false;
1710 }
1711 shader->info.uses_instanceid |= ls->info.uses_instanceid;
1712 parts[1] = ctx.main_fn;
1713
1714 /* LS prolog */
1715 if (vs_needs_prolog) {
1716 union si_shader_part_key vs_prolog_key;
1717 si_get_vs_prolog_key(&ls->info, shader_ls.info.num_input_sgprs, false,
1718 &shader->key.part.tcs.ls_prolog, shader, &vs_prolog_key);
1719 vs_prolog_key.vs_prolog.is_monolithic = true;
1720 si_llvm_build_vs_prolog(&ctx, &vs_prolog_key);
1721 parts[0] = ctx.main_fn;
1722 }
1723
1724 /* Reset the shader context. */
1725 ctx.shader = shader;
1726 ctx.type = PIPE_SHADER_TESS_CTRL;
1727
1728 si_build_wrapper_function(&ctx, parts + !vs_needs_prolog, 4 - !vs_needs_prolog,
1729 vs_needs_prolog, vs_needs_prolog ? 2 : 1);
1730 } else {
1731 LLVMValueRef parts[2];
1732 union si_shader_part_key epilog_key;
1733
1734 parts[0] = ctx.main_fn;
1735
1736 memset(&epilog_key, 0, sizeof(epilog_key));
1737 epilog_key.tcs_epilog.states = shader->key.part.tcs.epilog;
1738 si_llvm_build_tcs_epilog(&ctx, &epilog_key);
1739 parts[1] = ctx.main_fn;
1740
1741 si_build_wrapper_function(&ctx, parts, 2, 0, 0);
1742 }
1743 } else if (shader->is_monolithic && ctx.type == PIPE_SHADER_GEOMETRY) {
1744 if (ctx.screen->info.chip_class >= GFX9) {
1745 struct si_shader_selector *es = shader->key.part.gs.es;
1746 LLVMValueRef es_prolog = NULL;
1747 LLVMValueRef es_main = NULL;
1748 LLVMValueRef gs_prolog = NULL;
1749 LLVMValueRef gs_main = ctx.main_fn;
1750
1751 /* GS prolog */
1752 union si_shader_part_key gs_prolog_key;
1753 memset(&gs_prolog_key, 0, sizeof(gs_prolog_key));
1754 gs_prolog_key.gs_prolog.states = shader->key.part.gs.prolog;
1755 gs_prolog_key.gs_prolog.is_monolithic = true;
1756 gs_prolog_key.gs_prolog.as_ngg = shader->key.as_ngg;
1757 si_llvm_build_gs_prolog(&ctx, &gs_prolog_key);
1758 gs_prolog = ctx.main_fn;
1759
1760 /* ES main part */
1761 nir = get_nir_shader(es, &free_nir);
1762 struct si_shader shader_es = {};
1763 shader_es.selector = es;
1764 shader_es.key.as_es = 1;
1765 shader_es.key.as_ngg = shader->key.as_ngg;
1766 shader_es.key.mono = shader->key.mono;
1767 shader_es.key.opt = shader->key.opt;
1768 shader_es.is_monolithic = true;
1769
1770 if (!si_build_main_function(&ctx, &shader_es, nir, free_nir, false)) {
1771 si_llvm_dispose(&ctx);
1772 return false;
1773 }
1774 shader->info.uses_instanceid |= es->info.uses_instanceid;
1775 es_main = ctx.main_fn;
1776
1777 /* ES prolog */
1778 if (es->type == PIPE_SHADER_VERTEX &&
1779 si_vs_needs_prolog(es, &shader->key.part.gs.vs_prolog, &shader->key, false)) {
1780 union si_shader_part_key vs_prolog_key;
1781 si_get_vs_prolog_key(&es->info, shader_es.info.num_input_sgprs, false,
1782 &shader->key.part.gs.vs_prolog, shader, &vs_prolog_key);
1783 vs_prolog_key.vs_prolog.is_monolithic = true;
1784 si_llvm_build_vs_prolog(&ctx, &vs_prolog_key);
1785 es_prolog = ctx.main_fn;
1786 }
1787
1788 /* Reset the shader context. */
1789 ctx.shader = shader;
1790 ctx.type = PIPE_SHADER_GEOMETRY;
1791
1792 /* Prepare the array of shader parts. */
1793 LLVMValueRef parts[4];
1794 unsigned num_parts = 0, main_part, next_first_part;
1795
1796 if (es_prolog)
1797 parts[num_parts++] = es_prolog;
1798
1799 parts[main_part = num_parts++] = es_main;
1800 parts[next_first_part = num_parts++] = gs_prolog;
1801 parts[num_parts++] = gs_main;
1802
1803 si_build_wrapper_function(&ctx, parts, num_parts, main_part, next_first_part);
1804 } else {
1805 LLVMValueRef parts[2];
1806 union si_shader_part_key prolog_key;
1807
1808 parts[1] = ctx.main_fn;
1809
1810 memset(&prolog_key, 0, sizeof(prolog_key));
1811 prolog_key.gs_prolog.states = shader->key.part.gs.prolog;
1812 si_llvm_build_gs_prolog(&ctx, &prolog_key);
1813 parts[0] = ctx.main_fn;
1814
1815 si_build_wrapper_function(&ctx, parts, 2, 1, 0);
1816 }
1817 } else if (shader->is_monolithic && ctx.type == PIPE_SHADER_FRAGMENT) {
1818 si_llvm_build_monolithic_ps(&ctx, shader);
1819 }
1820
1821 si_llvm_optimize_module(&ctx);
1822
1823 /* Post-optimization transformations and analysis. */
1824 si_optimize_vs_outputs(&ctx);
1825
1826 if ((debug && debug->debug_message) || si_can_dump_shader(sscreen, ctx.type)) {
1827 ctx.shader->info.private_mem_vgprs = ac_count_scratch_private_memory(ctx.main_fn);
1828 }
1829
1830 /* Make sure the input is a pointer and not integer followed by inttoptr. */
1831 assert(LLVMGetTypeKind(LLVMTypeOf(LLVMGetParam(ctx.main_fn, 0))) == LLVMPointerTypeKind);
1832
1833 /* Compile to bytecode. */
1834 if (!si_compile_llvm(sscreen, &shader->binary, &shader->config, compiler, &ctx.ac, debug,
1835 ctx.type, si_get_shader_name(shader),
1836 si_should_optimize_less(compiler, shader->selector))) {
1837 si_llvm_dispose(&ctx);
1838 fprintf(stderr, "LLVM failed to compile shader\n");
1839 return false;
1840 }
1841
1842 si_llvm_dispose(&ctx);
1843 return true;
1844 }
1845
1846 bool si_compile_shader(struct si_screen *sscreen, struct ac_llvm_compiler *compiler,
1847 struct si_shader *shader, struct pipe_debug_callback *debug)
1848 {
1849 struct si_shader_selector *sel = shader->selector;
1850 bool free_nir;
1851 struct nir_shader *nir = get_nir_shader(sel, &free_nir);
1852
1853 /* Dump NIR before doing NIR->LLVM conversion in case the
1854 * conversion fails. */
1855 if (si_can_dump_shader(sscreen, sel->type) && !(sscreen->debug_flags & DBG(NO_NIR))) {
1856 nir_print_shader(nir, stderr);
1857 si_dump_streamout(&sel->so);
1858 }
1859
1860 memset(shader->info.vs_output_param_offset, AC_EXP_PARAM_UNDEFINED,
1861 sizeof(shader->info.vs_output_param_offset));
1862
1863 shader->info.uses_instanceid = sel->info.uses_instanceid;
1864
1865 /* TODO: ACO could compile non-monolithic shaders here (starting
1866 * with PS and NGG VS), but monolithic shaders should be compiled
1867 * by LLVM due to more complicated compilation.
1868 */
1869 if (!si_llvm_compile_shader(sscreen, compiler, shader, debug, nir, free_nir))
1870 return false;
1871
1872 /* Validate SGPR and VGPR usage for compute to detect compiler bugs.
1873 * LLVM 3.9svn has this bug.
1874 */
1875 if (sel->type == PIPE_SHADER_COMPUTE) {
1876 unsigned wave_size = sscreen->compute_wave_size;
1877 unsigned max_vgprs =
1878 sscreen->info.num_physical_wave64_vgprs_per_simd * (wave_size == 32 ? 2 : 1);
1879 unsigned max_sgprs = sscreen->info.num_physical_sgprs_per_simd;
1880 unsigned max_sgprs_per_wave = 128;
1881 unsigned simds_per_tg = 4; /* assuming WGP mode on gfx10 */
1882 unsigned threads_per_tg = si_get_max_workgroup_size(shader);
1883 unsigned waves_per_tg = DIV_ROUND_UP(threads_per_tg, wave_size);
1884 unsigned waves_per_simd = DIV_ROUND_UP(waves_per_tg, simds_per_tg);
1885
1886 max_vgprs = max_vgprs / waves_per_simd;
1887 max_sgprs = MIN2(max_sgprs / waves_per_simd, max_sgprs_per_wave);
1888
1889 if (shader->config.num_sgprs > max_sgprs || shader->config.num_vgprs > max_vgprs) {
1890 fprintf(stderr,
1891 "LLVM failed to compile a shader correctly: "
1892 "SGPR:VGPR usage is %u:%u, but the hw limit is %u:%u\n",
1893 shader->config.num_sgprs, shader->config.num_vgprs, max_sgprs, max_vgprs);
1894
1895 /* Just terminate the process, because dependent
1896 * shaders can hang due to bad input data, but use
1897 * the env var to allow shader-db to work.
1898 */
1899 if (!debug_get_bool_option("SI_PASS_BAD_SHADERS", false))
1900 abort();
1901 }
1902 }
1903
1904 /* Add the scratch offset to input SGPRs. */
1905 if (shader->config.scratch_bytes_per_wave && !si_is_merged_shader(shader))
1906 shader->info.num_input_sgprs += 1; /* scratch byte offset */
1907
1908 /* Calculate the number of fragment input VGPRs. */
1909 if (sel->type == PIPE_SHADER_FRAGMENT) {
1910 shader->info.num_input_vgprs = ac_get_fs_input_vgpr_cnt(
1911 &shader->config, &shader->info.face_vgpr_index, &shader->info.ancillary_vgpr_index);
1912 }
1913
1914 si_calculate_max_simd_waves(shader);
1915 si_shader_dump_stats_for_shader_db(sscreen, shader, debug);
1916 return true;
1917 }
1918
1919 /**
1920 * Create, compile and return a shader part (prolog or epilog).
1921 *
1922 * \param sscreen screen
1923 * \param list list of shader parts of the same category
1924 * \param type shader type
1925 * \param key shader part key
1926 * \param prolog whether the part being requested is a prolog
1927 * \param tm LLVM target machine
1928 * \param debug debug callback
1929 * \param build the callback responsible for building the main function
1930 * \return non-NULL on success
1931 */
1932 static struct si_shader_part *
1933 si_get_shader_part(struct si_screen *sscreen, struct si_shader_part **list,
1934 enum pipe_shader_type type, bool prolog, union si_shader_part_key *key,
1935 struct ac_llvm_compiler *compiler, struct pipe_debug_callback *debug,
1936 void (*build)(struct si_shader_context *, union si_shader_part_key *),
1937 const char *name)
1938 {
1939 struct si_shader_part *result;
1940
1941 simple_mtx_lock(&sscreen->shader_parts_mutex);
1942
1943 /* Find existing. */
1944 for (result = *list; result; result = result->next) {
1945 if (memcmp(&result->key, key, sizeof(*key)) == 0) {
1946 simple_mtx_unlock(&sscreen->shader_parts_mutex);
1947 return result;
1948 }
1949 }
1950
1951 /* Compile a new one. */
1952 result = CALLOC_STRUCT(si_shader_part);
1953 result->key = *key;
1954
1955 struct si_shader_selector sel = {};
1956 sel.screen = sscreen;
1957
1958 struct si_shader shader = {};
1959 shader.selector = &sel;
1960
1961 switch (type) {
1962 case PIPE_SHADER_VERTEX:
1963 shader.key.as_ls = key->vs_prolog.as_ls;
1964 shader.key.as_es = key->vs_prolog.as_es;
1965 shader.key.as_ngg = key->vs_prolog.as_ngg;
1966 shader.key.opt.ngg_culling =
1967 (key->vs_prolog.gs_fast_launch_tri_list ? SI_NGG_CULL_GS_FAST_LAUNCH_TRI_LIST : 0) |
1968 (key->vs_prolog.gs_fast_launch_tri_strip ? SI_NGG_CULL_GS_FAST_LAUNCH_TRI_STRIP : 0);
1969 shader.key.opt.vs_as_prim_discard_cs = key->vs_prolog.as_prim_discard_cs;
1970 break;
1971 case PIPE_SHADER_TESS_CTRL:
1972 assert(!prolog);
1973 shader.key.part.tcs.epilog = key->tcs_epilog.states;
1974 break;
1975 case PIPE_SHADER_GEOMETRY:
1976 assert(prolog);
1977 shader.key.as_ngg = key->gs_prolog.as_ngg;
1978 break;
1979 case PIPE_SHADER_FRAGMENT:
1980 if (prolog)
1981 shader.key.part.ps.prolog = key->ps_prolog.states;
1982 else
1983 shader.key.part.ps.epilog = key->ps_epilog.states;
1984 break;
1985 default:
1986 unreachable("bad shader part");
1987 }
1988
1989 struct si_shader_context ctx;
1990 si_llvm_context_init(&ctx, sscreen, compiler,
1991 si_get_wave_size(sscreen, type, shader.key.as_ngg, shader.key.as_es,
1992 shader.key.opt.ngg_culling & SI_NGG_CULL_GS_FAST_LAUNCH_ALL,
1993 shader.key.opt.vs_as_prim_discard_cs));
1994 ctx.shader = &shader;
1995 ctx.type = type;
1996
1997 build(&ctx, key);
1998
1999 /* Compile. */
2000 si_llvm_optimize_module(&ctx);
2001
2002 if (!si_compile_llvm(sscreen, &result->binary, &result->config, compiler, &ctx.ac, debug,
2003 ctx.type, name, false)) {
2004 FREE(result);
2005 result = NULL;
2006 goto out;
2007 }
2008
2009 result->next = *list;
2010 *list = result;
2011
2012 out:
2013 si_llvm_dispose(&ctx);
2014 simple_mtx_unlock(&sscreen->shader_parts_mutex);
2015 return result;
2016 }
2017
2018 static bool si_get_vs_prolog(struct si_screen *sscreen, struct ac_llvm_compiler *compiler,
2019 struct si_shader *shader, struct pipe_debug_callback *debug,
2020 struct si_shader *main_part, const struct si_vs_prolog_bits *key)
2021 {
2022 struct si_shader_selector *vs = main_part->selector;
2023
2024 if (!si_vs_needs_prolog(vs, key, &shader->key, false))
2025 return true;
2026
2027 /* Get the prolog. */
2028 union si_shader_part_key prolog_key;
2029 si_get_vs_prolog_key(&vs->info, main_part->info.num_input_sgprs, false, key, shader,
2030 &prolog_key);
2031
2032 shader->prolog =
2033 si_get_shader_part(sscreen, &sscreen->vs_prologs, PIPE_SHADER_VERTEX, true, &prolog_key,
2034 compiler, debug, si_llvm_build_vs_prolog, "Vertex Shader Prolog");
2035 return shader->prolog != NULL;
2036 }
2037
2038 /**
2039 * Select and compile (or reuse) vertex shader parts (prolog & epilog).
2040 */
2041 static bool si_shader_select_vs_parts(struct si_screen *sscreen, struct ac_llvm_compiler *compiler,
2042 struct si_shader *shader, struct pipe_debug_callback *debug)
2043 {
2044 return si_get_vs_prolog(sscreen, compiler, shader, debug, shader, &shader->key.part.vs.prolog);
2045 }
2046
2047 /**
2048 * Select and compile (or reuse) TCS parts (epilog).
2049 */
2050 static bool si_shader_select_tcs_parts(struct si_screen *sscreen, struct ac_llvm_compiler *compiler,
2051 struct si_shader *shader, struct pipe_debug_callback *debug)
2052 {
2053 if (sscreen->info.chip_class >= GFX9) {
2054 struct si_shader *ls_main_part = shader->key.part.tcs.ls->main_shader_part_ls;
2055
2056 if (!si_get_vs_prolog(sscreen, compiler, shader, debug, ls_main_part,
2057 &shader->key.part.tcs.ls_prolog))
2058 return false;
2059
2060 shader->previous_stage = ls_main_part;
2061 }
2062
2063 /* Get the epilog. */
2064 union si_shader_part_key epilog_key;
2065 memset(&epilog_key, 0, sizeof(epilog_key));
2066 epilog_key.tcs_epilog.states = shader->key.part.tcs.epilog;
2067
2068 shader->epilog = si_get_shader_part(sscreen, &sscreen->tcs_epilogs, PIPE_SHADER_TESS_CTRL, false,
2069 &epilog_key, compiler, debug, si_llvm_build_tcs_epilog,
2070 "Tessellation Control Shader Epilog");
2071 return shader->epilog != NULL;
2072 }
2073
2074 /**
2075 * Select and compile (or reuse) GS parts (prolog).
2076 */
2077 static bool si_shader_select_gs_parts(struct si_screen *sscreen, struct ac_llvm_compiler *compiler,
2078 struct si_shader *shader, struct pipe_debug_callback *debug)
2079 {
2080 if (sscreen->info.chip_class >= GFX9) {
2081 struct si_shader *es_main_part;
2082 enum pipe_shader_type es_type = shader->key.part.gs.es->type;
2083
2084 if (shader->key.as_ngg)
2085 es_main_part = shader->key.part.gs.es->main_shader_part_ngg_es;
2086 else
2087 es_main_part = shader->key.part.gs.es->main_shader_part_es;
2088
2089 if (es_type == PIPE_SHADER_VERTEX &&
2090 !si_get_vs_prolog(sscreen, compiler, shader, debug, es_main_part,
2091 &shader->key.part.gs.vs_prolog))
2092 return false;
2093
2094 shader->previous_stage = es_main_part;
2095 }
2096
2097 if (!shader->key.part.gs.prolog.tri_strip_adj_fix)
2098 return true;
2099
2100 union si_shader_part_key prolog_key;
2101 memset(&prolog_key, 0, sizeof(prolog_key));
2102 prolog_key.gs_prolog.states = shader->key.part.gs.prolog;
2103 prolog_key.gs_prolog.as_ngg = shader->key.as_ngg;
2104
2105 shader->prolog2 =
2106 si_get_shader_part(sscreen, &sscreen->gs_prologs, PIPE_SHADER_GEOMETRY, true, &prolog_key,
2107 compiler, debug, si_llvm_build_gs_prolog, "Geometry Shader Prolog");
2108 return shader->prolog2 != NULL;
2109 }
2110
2111 /**
2112 * Compute the PS prolog key, which contains all the information needed to
2113 * build the PS prolog function, and set related bits in shader->config.
2114 */
2115 void si_get_ps_prolog_key(struct si_shader *shader, union si_shader_part_key *key,
2116 bool separate_prolog)
2117 {
2118 struct si_shader_info *info = &shader->selector->info;
2119
2120 memset(key, 0, sizeof(*key));
2121 key->ps_prolog.states = shader->key.part.ps.prolog;
2122 key->ps_prolog.colors_read = info->colors_read;
2123 key->ps_prolog.num_input_sgprs = shader->info.num_input_sgprs;
2124 key->ps_prolog.num_input_vgprs = shader->info.num_input_vgprs;
2125 key->ps_prolog.wqm =
2126 info->uses_derivatives &&
2127 (key->ps_prolog.colors_read || key->ps_prolog.states.force_persp_sample_interp ||
2128 key->ps_prolog.states.force_linear_sample_interp ||
2129 key->ps_prolog.states.force_persp_center_interp ||
2130 key->ps_prolog.states.force_linear_center_interp ||
2131 key->ps_prolog.states.bc_optimize_for_persp || key->ps_prolog.states.bc_optimize_for_linear);
2132 key->ps_prolog.ancillary_vgpr_index = shader->info.ancillary_vgpr_index;
2133
2134 if (info->colors_read) {
2135 unsigned *color = shader->selector->color_attr_index;
2136
2137 if (shader->key.part.ps.prolog.color_two_side) {
2138 /* BCOLORs are stored after the last input. */
2139 key->ps_prolog.num_interp_inputs = info->num_inputs;
2140 key->ps_prolog.face_vgpr_index = shader->info.face_vgpr_index;
2141 if (separate_prolog)
2142 shader->config.spi_ps_input_ena |= S_0286CC_FRONT_FACE_ENA(1);
2143 }
2144
2145 for (unsigned i = 0; i < 2; i++) {
2146 unsigned interp = info->input_interpolate[color[i]];
2147 unsigned location = info->input_interpolate_loc[color[i]];
2148
2149 if (!(info->colors_read & (0xf << i * 4)))
2150 continue;
2151
2152 key->ps_prolog.color_attr_index[i] = color[i];
2153
2154 if (shader->key.part.ps.prolog.flatshade_colors && interp == TGSI_INTERPOLATE_COLOR)
2155 interp = TGSI_INTERPOLATE_CONSTANT;
2156
2157 switch (interp) {
2158 case TGSI_INTERPOLATE_CONSTANT:
2159 key->ps_prolog.color_interp_vgpr_index[i] = -1;
2160 break;
2161 case TGSI_INTERPOLATE_PERSPECTIVE:
2162 case TGSI_INTERPOLATE_COLOR:
2163 /* Force the interpolation location for colors here. */
2164 if (shader->key.part.ps.prolog.force_persp_sample_interp)
2165 location = TGSI_INTERPOLATE_LOC_SAMPLE;
2166 if (shader->key.part.ps.prolog.force_persp_center_interp)
2167 location = TGSI_INTERPOLATE_LOC_CENTER;
2168
2169 switch (location) {
2170 case TGSI_INTERPOLATE_LOC_SAMPLE:
2171 key->ps_prolog.color_interp_vgpr_index[i] = 0;
2172 if (separate_prolog) {
2173 shader->config.spi_ps_input_ena |= S_0286CC_PERSP_SAMPLE_ENA(1);
2174 }
2175 break;
2176 case TGSI_INTERPOLATE_LOC_CENTER:
2177 key->ps_prolog.color_interp_vgpr_index[i] = 2;
2178 if (separate_prolog) {
2179 shader->config.spi_ps_input_ena |= S_0286CC_PERSP_CENTER_ENA(1);
2180 }
2181 break;
2182 case TGSI_INTERPOLATE_LOC_CENTROID:
2183 key->ps_prolog.color_interp_vgpr_index[i] = 4;
2184 if (separate_prolog) {
2185 shader->config.spi_ps_input_ena |= S_0286CC_PERSP_CENTROID_ENA(1);
2186 }
2187 break;
2188 default:
2189 assert(0);
2190 }
2191 break;
2192 case TGSI_INTERPOLATE_LINEAR:
2193 /* Force the interpolation location for colors here. */
2194 if (shader->key.part.ps.prolog.force_linear_sample_interp)
2195 location = TGSI_INTERPOLATE_LOC_SAMPLE;
2196 if (shader->key.part.ps.prolog.force_linear_center_interp)
2197 location = TGSI_INTERPOLATE_LOC_CENTER;
2198
2199 /* The VGPR assignment for non-monolithic shaders
2200 * works because InitialPSInputAddr is set on the
2201 * main shader and PERSP_PULL_MODEL is never used.
2202 */
2203 switch (location) {
2204 case TGSI_INTERPOLATE_LOC_SAMPLE:
2205 key->ps_prolog.color_interp_vgpr_index[i] = separate_prolog ? 6 : 9;
2206 if (separate_prolog) {
2207 shader->config.spi_ps_input_ena |= S_0286CC_LINEAR_SAMPLE_ENA(1);
2208 }
2209 break;
2210 case TGSI_INTERPOLATE_LOC_CENTER:
2211 key->ps_prolog.color_interp_vgpr_index[i] = separate_prolog ? 8 : 11;
2212 if (separate_prolog) {
2213 shader->config.spi_ps_input_ena |= S_0286CC_LINEAR_CENTER_ENA(1);
2214 }
2215 break;
2216 case TGSI_INTERPOLATE_LOC_CENTROID:
2217 key->ps_prolog.color_interp_vgpr_index[i] = separate_prolog ? 10 : 13;
2218 if (separate_prolog) {
2219 shader->config.spi_ps_input_ena |= S_0286CC_LINEAR_CENTROID_ENA(1);
2220 }
2221 break;
2222 default:
2223 assert(0);
2224 }
2225 break;
2226 default:
2227 assert(0);
2228 }
2229 }
2230 }
2231 }
2232
2233 /**
2234 * Check whether a PS prolog is required based on the key.
2235 */
2236 bool si_need_ps_prolog(const union si_shader_part_key *key)
2237 {
2238 return key->ps_prolog.colors_read || key->ps_prolog.states.force_persp_sample_interp ||
2239 key->ps_prolog.states.force_linear_sample_interp ||
2240 key->ps_prolog.states.force_persp_center_interp ||
2241 key->ps_prolog.states.force_linear_center_interp ||
2242 key->ps_prolog.states.bc_optimize_for_persp ||
2243 key->ps_prolog.states.bc_optimize_for_linear || key->ps_prolog.states.poly_stipple ||
2244 key->ps_prolog.states.samplemask_log_ps_iter;
2245 }
2246
2247 /**
2248 * Compute the PS epilog key, which contains all the information needed to
2249 * build the PS epilog function.
2250 */
2251 void si_get_ps_epilog_key(struct si_shader *shader, union si_shader_part_key *key)
2252 {
2253 struct si_shader_info *info = &shader->selector->info;
2254 memset(key, 0, sizeof(*key));
2255 key->ps_epilog.colors_written = info->colors_written;
2256 key->ps_epilog.writes_z = info->writes_z;
2257 key->ps_epilog.writes_stencil = info->writes_stencil;
2258 key->ps_epilog.writes_samplemask = info->writes_samplemask;
2259 key->ps_epilog.states = shader->key.part.ps.epilog;
2260 }
2261
2262 /**
2263 * Select and compile (or reuse) pixel shader parts (prolog & epilog).
2264 */
2265 static bool si_shader_select_ps_parts(struct si_screen *sscreen, struct ac_llvm_compiler *compiler,
2266 struct si_shader *shader, struct pipe_debug_callback *debug)
2267 {
2268 union si_shader_part_key prolog_key;
2269 union si_shader_part_key epilog_key;
2270
2271 /* Get the prolog. */
2272 si_get_ps_prolog_key(shader, &prolog_key, true);
2273
2274 /* The prolog is a no-op if these aren't set. */
2275 if (si_need_ps_prolog(&prolog_key)) {
2276 shader->prolog =
2277 si_get_shader_part(sscreen, &sscreen->ps_prologs, PIPE_SHADER_FRAGMENT, true, &prolog_key,
2278 compiler, debug, si_llvm_build_ps_prolog, "Fragment Shader Prolog");
2279 if (!shader->prolog)
2280 return false;
2281 }
2282
2283 /* Get the epilog. */
2284 si_get_ps_epilog_key(shader, &epilog_key);
2285
2286 shader->epilog =
2287 si_get_shader_part(sscreen, &sscreen->ps_epilogs, PIPE_SHADER_FRAGMENT, false, &epilog_key,
2288 compiler, debug, si_llvm_build_ps_epilog, "Fragment Shader Epilog");
2289 if (!shader->epilog)
2290 return false;
2291
2292 /* Enable POS_FIXED_PT if polygon stippling is enabled. */
2293 if (shader->key.part.ps.prolog.poly_stipple) {
2294 shader->config.spi_ps_input_ena |= S_0286CC_POS_FIXED_PT_ENA(1);
2295 assert(G_0286CC_POS_FIXED_PT_ENA(shader->config.spi_ps_input_addr));
2296 }
2297
2298 /* Set up the enable bits for per-sample shading if needed. */
2299 if (shader->key.part.ps.prolog.force_persp_sample_interp &&
2300 (G_0286CC_PERSP_CENTER_ENA(shader->config.spi_ps_input_ena) ||
2301 G_0286CC_PERSP_CENTROID_ENA(shader->config.spi_ps_input_ena))) {
2302 shader->config.spi_ps_input_ena &= C_0286CC_PERSP_CENTER_ENA;
2303 shader->config.spi_ps_input_ena &= C_0286CC_PERSP_CENTROID_ENA;
2304 shader->config.spi_ps_input_ena |= S_0286CC_PERSP_SAMPLE_ENA(1);
2305 }
2306 if (shader->key.part.ps.prolog.force_linear_sample_interp &&
2307 (G_0286CC_LINEAR_CENTER_ENA(shader->config.spi_ps_input_ena) ||
2308 G_0286CC_LINEAR_CENTROID_ENA(shader->config.spi_ps_input_ena))) {
2309 shader->config.spi_ps_input_ena &= C_0286CC_LINEAR_CENTER_ENA;
2310 shader->config.spi_ps_input_ena &= C_0286CC_LINEAR_CENTROID_ENA;
2311 shader->config.spi_ps_input_ena |= S_0286CC_LINEAR_SAMPLE_ENA(1);
2312 }
2313 if (shader->key.part.ps.prolog.force_persp_center_interp &&
2314 (G_0286CC_PERSP_SAMPLE_ENA(shader->config.spi_ps_input_ena) ||
2315 G_0286CC_PERSP_CENTROID_ENA(shader->config.spi_ps_input_ena))) {
2316 shader->config.spi_ps_input_ena &= C_0286CC_PERSP_SAMPLE_ENA;
2317 shader->config.spi_ps_input_ena &= C_0286CC_PERSP_CENTROID_ENA;
2318 shader->config.spi_ps_input_ena |= S_0286CC_PERSP_CENTER_ENA(1);
2319 }
2320 if (shader->key.part.ps.prolog.force_linear_center_interp &&
2321 (G_0286CC_LINEAR_SAMPLE_ENA(shader->config.spi_ps_input_ena) ||
2322 G_0286CC_LINEAR_CENTROID_ENA(shader->config.spi_ps_input_ena))) {
2323 shader->config.spi_ps_input_ena &= C_0286CC_LINEAR_SAMPLE_ENA;
2324 shader->config.spi_ps_input_ena &= C_0286CC_LINEAR_CENTROID_ENA;
2325 shader->config.spi_ps_input_ena |= S_0286CC_LINEAR_CENTER_ENA(1);
2326 }
2327
2328 /* POW_W_FLOAT requires that one of the perspective weights is enabled. */
2329 if (G_0286CC_POS_W_FLOAT_ENA(shader->config.spi_ps_input_ena) &&
2330 !(shader->config.spi_ps_input_ena & 0xf)) {
2331 shader->config.spi_ps_input_ena |= S_0286CC_PERSP_CENTER_ENA(1);
2332 assert(G_0286CC_PERSP_CENTER_ENA(shader->config.spi_ps_input_addr));
2333 }
2334
2335 /* At least one pair of interpolation weights must be enabled. */
2336 if (!(shader->config.spi_ps_input_ena & 0x7f)) {
2337 shader->config.spi_ps_input_ena |= S_0286CC_LINEAR_CENTER_ENA(1);
2338 assert(G_0286CC_LINEAR_CENTER_ENA(shader->config.spi_ps_input_addr));
2339 }
2340
2341 /* Samplemask fixup requires the sample ID. */
2342 if (shader->key.part.ps.prolog.samplemask_log_ps_iter) {
2343 shader->config.spi_ps_input_ena |= S_0286CC_ANCILLARY_ENA(1);
2344 assert(G_0286CC_ANCILLARY_ENA(shader->config.spi_ps_input_addr));
2345 }
2346
2347 /* The sample mask input is always enabled, because the API shader always
2348 * passes it through to the epilog. Disable it here if it's unused.
2349 */
2350 if (!shader->key.part.ps.epilog.poly_line_smoothing && !shader->selector->info.reads_samplemask)
2351 shader->config.spi_ps_input_ena &= C_0286CC_SAMPLE_COVERAGE_ENA;
2352
2353 return true;
2354 }
2355
2356 void si_multiwave_lds_size_workaround(struct si_screen *sscreen, unsigned *lds_size)
2357 {
2358 /* If tessellation is all offchip and on-chip GS isn't used, this
2359 * workaround is not needed.
2360 */
2361 return;
2362
2363 /* SPI barrier management bug:
2364 * Make sure we have at least 4k of LDS in use to avoid the bug.
2365 * It applies to workgroup sizes of more than one wavefront.
2366 */
2367 if (sscreen->info.family == CHIP_BONAIRE || sscreen->info.family == CHIP_KABINI)
2368 *lds_size = MAX2(*lds_size, 8);
2369 }
2370
2371 void si_fix_resource_usage(struct si_screen *sscreen, struct si_shader *shader)
2372 {
2373 unsigned min_sgprs = shader->info.num_input_sgprs + 2; /* VCC */
2374
2375 shader->config.num_sgprs = MAX2(shader->config.num_sgprs, min_sgprs);
2376
2377 if (shader->selector->type == PIPE_SHADER_COMPUTE &&
2378 si_get_max_workgroup_size(shader) > sscreen->compute_wave_size) {
2379 si_multiwave_lds_size_workaround(sscreen, &shader->config.lds_size);
2380 }
2381 }
2382
2383 bool si_create_shader_variant(struct si_screen *sscreen, struct ac_llvm_compiler *compiler,
2384 struct si_shader *shader, struct pipe_debug_callback *debug)
2385 {
2386 struct si_shader_selector *sel = shader->selector;
2387 struct si_shader *mainp = *si_get_main_shader_part(sel, &shader->key);
2388
2389 /* LS, ES, VS are compiled on demand if the main part hasn't been
2390 * compiled for that stage.
2391 *
2392 * GS are compiled on demand if the main part hasn't been compiled
2393 * for the chosen NGG-ness.
2394 *
2395 * Vertex shaders are compiled on demand when a vertex fetch
2396 * workaround must be applied.
2397 */
2398 if (shader->is_monolithic) {
2399 /* Monolithic shader (compiled as a whole, has many variants,
2400 * may take a long time to compile).
2401 */
2402 if (!si_compile_shader(sscreen, compiler, shader, debug))
2403 return false;
2404 } else {
2405 /* The shader consists of several parts:
2406 *
2407 * - the middle part is the user shader, it has 1 variant only
2408 * and it was compiled during the creation of the shader
2409 * selector
2410 * - the prolog part is inserted at the beginning
2411 * - the epilog part is inserted at the end
2412 *
2413 * The prolog and epilog have many (but simple) variants.
2414 *
2415 * Starting with gfx9, geometry and tessellation control
2416 * shaders also contain the prolog and user shader parts of
2417 * the previous shader stage.
2418 */
2419
2420 if (!mainp)
2421 return false;
2422
2423 /* Copy the compiled shader data over. */
2424 shader->is_binary_shared = true;
2425 shader->binary = mainp->binary;
2426 shader->config = mainp->config;
2427 shader->info.num_input_sgprs = mainp->info.num_input_sgprs;
2428 shader->info.num_input_vgprs = mainp->info.num_input_vgprs;
2429 shader->info.face_vgpr_index = mainp->info.face_vgpr_index;
2430 shader->info.ancillary_vgpr_index = mainp->info.ancillary_vgpr_index;
2431 memcpy(shader->info.vs_output_param_offset, mainp->info.vs_output_param_offset,
2432 sizeof(mainp->info.vs_output_param_offset));
2433 shader->info.uses_instanceid = mainp->info.uses_instanceid;
2434 shader->info.nr_pos_exports = mainp->info.nr_pos_exports;
2435 shader->info.nr_param_exports = mainp->info.nr_param_exports;
2436
2437 /* Select prologs and/or epilogs. */
2438 switch (sel->type) {
2439 case PIPE_SHADER_VERTEX:
2440 if (!si_shader_select_vs_parts(sscreen, compiler, shader, debug))
2441 return false;
2442 break;
2443 case PIPE_SHADER_TESS_CTRL:
2444 if (!si_shader_select_tcs_parts(sscreen, compiler, shader, debug))
2445 return false;
2446 break;
2447 case PIPE_SHADER_TESS_EVAL:
2448 break;
2449 case PIPE_SHADER_GEOMETRY:
2450 if (!si_shader_select_gs_parts(sscreen, compiler, shader, debug))
2451 return false;
2452 break;
2453 case PIPE_SHADER_FRAGMENT:
2454 if (!si_shader_select_ps_parts(sscreen, compiler, shader, debug))
2455 return false;
2456
2457 /* Make sure we have at least as many VGPRs as there
2458 * are allocated inputs.
2459 */
2460 shader->config.num_vgprs = MAX2(shader->config.num_vgprs, shader->info.num_input_vgprs);
2461 break;
2462 default:;
2463 }
2464
2465 /* Update SGPR and VGPR counts. */
2466 if (shader->prolog) {
2467 shader->config.num_sgprs =
2468 MAX2(shader->config.num_sgprs, shader->prolog->config.num_sgprs);
2469 shader->config.num_vgprs =
2470 MAX2(shader->config.num_vgprs, shader->prolog->config.num_vgprs);
2471 }
2472 if (shader->previous_stage) {
2473 shader->config.num_sgprs =
2474 MAX2(shader->config.num_sgprs, shader->previous_stage->config.num_sgprs);
2475 shader->config.num_vgprs =
2476 MAX2(shader->config.num_vgprs, shader->previous_stage->config.num_vgprs);
2477 shader->config.spilled_sgprs =
2478 MAX2(shader->config.spilled_sgprs, shader->previous_stage->config.spilled_sgprs);
2479 shader->config.spilled_vgprs =
2480 MAX2(shader->config.spilled_vgprs, shader->previous_stage->config.spilled_vgprs);
2481 shader->info.private_mem_vgprs =
2482 MAX2(shader->info.private_mem_vgprs, shader->previous_stage->info.private_mem_vgprs);
2483 shader->config.scratch_bytes_per_wave =
2484 MAX2(shader->config.scratch_bytes_per_wave,
2485 shader->previous_stage->config.scratch_bytes_per_wave);
2486 shader->info.uses_instanceid |= shader->previous_stage->info.uses_instanceid;
2487 }
2488 if (shader->prolog2) {
2489 shader->config.num_sgprs =
2490 MAX2(shader->config.num_sgprs, shader->prolog2->config.num_sgprs);
2491 shader->config.num_vgprs =
2492 MAX2(shader->config.num_vgprs, shader->prolog2->config.num_vgprs);
2493 }
2494 if (shader->epilog) {
2495 shader->config.num_sgprs =
2496 MAX2(shader->config.num_sgprs, shader->epilog->config.num_sgprs);
2497 shader->config.num_vgprs =
2498 MAX2(shader->config.num_vgprs, shader->epilog->config.num_vgprs);
2499 }
2500 si_calculate_max_simd_waves(shader);
2501 }
2502
2503 if (shader->key.as_ngg) {
2504 assert(!shader->key.as_es && !shader->key.as_ls);
2505 if (!gfx10_ngg_calculate_subgroup_info(shader)) {
2506 fprintf(stderr, "Failed to compute subgroup info\n");
2507 return false;
2508 }
2509 } else if (sscreen->info.chip_class >= GFX9 && sel->type == PIPE_SHADER_GEOMETRY) {
2510 gfx9_get_gs_info(shader->previous_stage_sel, sel, &shader->gs_info);
2511 }
2512
2513 si_fix_resource_usage(sscreen, shader);
2514 si_shader_dump(sscreen, shader, debug, stderr, true);
2515
2516 /* Upload. */
2517 if (!si_shader_binary_upload(sscreen, shader, 0)) {
2518 fprintf(stderr, "LLVM failed to upload shader\n");
2519 return false;
2520 }
2521
2522 return true;
2523 }
2524
2525 void si_shader_binary_clean(struct si_shader_binary *binary)
2526 {
2527 free((void *)binary->elf_buffer);
2528 binary->elf_buffer = NULL;
2529
2530 free(binary->llvm_ir_string);
2531 binary->llvm_ir_string = NULL;
2532 }
2533
2534 void si_shader_destroy(struct si_shader *shader)
2535 {
2536 if (shader->scratch_bo)
2537 si_resource_reference(&shader->scratch_bo, NULL);
2538
2539 si_resource_reference(&shader->bo, NULL);
2540
2541 if (!shader->is_binary_shared)
2542 si_shader_binary_clean(&shader->binary);
2543
2544 free(shader->shader_log);
2545 }