projects / mesa.git / blob
? search:
summary | shortlog | log | commit | commitdiff | tree
history | raw | HEAD
radv: move setting can_discard to ac_fill_shader_info()
[mesa.git] / src / amd / vulkan / radv_nir_to_llvm.c
1 /*
2 * Copyright © 2016 Red Hat.
3 * Copyright © 2016 Bas Nieuwenhuizen
4 *
5 * based in part on anv driver which is:
6 * Copyright © 2015 Intel Corporation
7 *
8 * Permission is hereby granted, free of charge, to any person obtaining a
9 * copy of this software and associated documentation files (the "Software"),
10 * to deal in the Software without restriction, including without limitation
11 * the rights to use, copy, modify, merge, publish, distribute, sublicense,
12 * and/or sell copies of the Software, and to permit persons to whom the
13 * Software is furnished to do so, subject to the following conditions:
14 *
15 * The above copyright notice and this permission notice (including the next
16 * paragraph) shall be included in all copies or substantial portions of the
17 * Software.
18 *
19 * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
20 * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
21 * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL
22 * THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
23 * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING
24 * FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS
25 * IN THE SOFTWARE.
26 */
27
28 #include "radv_private.h"
29 #include "radv_shader.h"
30 #include "radv_shader_helper.h"
31 #include "nir/nir.h"
32
33 #include <llvm-c/Core.h>
34 #include <llvm-c/TargetMachine.h>
35 #include <llvm-c/Transforms/Scalar.h>
36 #include <llvm-c/Transforms/Utils.h>
37
38 #include "sid.h"
39 #include "ac_binary.h"
40 #include "ac_llvm_util.h"
41 #include "ac_llvm_build.h"
42 #include "ac_shader_abi.h"
43 #include "ac_shader_util.h"
44 #include "ac_exp_param.h"
45
46 #define RADEON_LLVM_MAX_INPUTS (VARYING_SLOT_VAR31 + 1)
47
48 struct radv_shader_context {
49 struct ac_llvm_context ac;
50 const struct radv_nir_compiler_options *options;
51 struct radv_shader_variant_info *shader_info;
52 struct ac_shader_abi abi;
53
54 unsigned max_workgroup_size;
55 LLVMContextRef context;
56 LLVMValueRef main_function;
57
58 LLVMValueRef descriptor_sets[MAX_SETS];
59 LLVMValueRef ring_offsets;
60
61 LLVMValueRef vertex_buffers;
62 LLVMValueRef rel_auto_id;
63 LLVMValueRef vs_prim_id;
64 LLVMValueRef es2gs_offset;
65
66 LLVMValueRef oc_lds;
67 LLVMValueRef merged_wave_info;
68 LLVMValueRef tess_factor_offset;
69 LLVMValueRef tes_rel_patch_id;
70 LLVMValueRef tes_u;
71 LLVMValueRef tes_v;
72
73 /* HW GS */
74 /* On gfx10:
75 * - bits 0..10: ordered_wave_id
76 * - bits 12..20: number of vertices in group
77 * - bits 22..30: number of primitives in group
78 */
79 LLVMValueRef gs_tg_info;
80 LLVMValueRef gs2vs_offset;
81 LLVMValueRef gs_wave_id;
82 LLVMValueRef gs_vtx_offset[6];
83
84 LLVMValueRef esgs_ring;
85 LLVMValueRef gsvs_ring[4];
86 LLVMValueRef hs_ring_tess_offchip;
87 LLVMValueRef hs_ring_tess_factor;
88
89 /* Streamout */
90 LLVMValueRef streamout_buffers;
91 LLVMValueRef streamout_write_idx;
92 LLVMValueRef streamout_config;
93 LLVMValueRef streamout_offset[4];
94
95 gl_shader_stage stage;
96
97 LLVMValueRef inputs[RADEON_LLVM_MAX_INPUTS * 4];
98
99 uint64_t output_mask;
100
101 bool is_gs_copy_shader;
102 LLVMValueRef gs_next_vertex[4];
103 LLVMValueRef gs_curprim_verts[4];
104 LLVMValueRef gs_generated_prims[4];
105 LLVMValueRef gs_ngg_emit;
106 LLVMValueRef gs_ngg_scratch;
107 unsigned gs_max_out_vertices;
108 unsigned gs_output_prim;
109
110 unsigned tes_primitive_mode;
111
112 uint32_t tcs_patch_outputs_read;
113 uint64_t tcs_outputs_read;
114 uint32_t tcs_vertices_per_patch;
115 uint32_t tcs_num_inputs;
116 uint32_t tcs_num_patches;
117 uint32_t max_gsvs_emit_size;
118 uint32_t gsvs_vertex_size;
119
120 LLVMValueRef vertexptr; /* GFX10 only */
121 };
122
123 struct radv_shader_output_values {
124 LLVMValueRef values[4];
125 unsigned slot_name;
126 unsigned slot_index;
127 unsigned usage_mask;
128 };
129
130 enum radeon_llvm_calling_convention {
131 RADEON_LLVM_AMDGPU_VS = 87,
132 RADEON_LLVM_AMDGPU_GS = 88,
133 RADEON_LLVM_AMDGPU_PS = 89,
134 RADEON_LLVM_AMDGPU_CS = 90,
135 RADEON_LLVM_AMDGPU_HS = 93,
136 };
137
138 static inline struct radv_shader_context *
139 radv_shader_context_from_abi(struct ac_shader_abi *abi)
140 {
141 struct radv_shader_context *ctx = NULL;
142 return container_of(abi, ctx, abi);
143 }
144
145 static LLVMValueRef get_rel_patch_id(struct radv_shader_context *ctx)
146 {
147 switch (ctx->stage) {
148 case MESA_SHADER_TESS_CTRL:
149 return ac_unpack_param(&ctx->ac, ctx->abi.tcs_rel_ids, 0, 8);
150 case MESA_SHADER_TESS_EVAL:
151 return ctx->tes_rel_patch_id;
152 break;
153 default:
154 unreachable("Illegal stage");
155 }
156 }
157
158 static unsigned
159 get_tcs_num_patches(struct radv_shader_context *ctx)
160 {
161 unsigned num_tcs_input_cp = ctx->options->key.tcs.input_vertices;
162 unsigned num_tcs_output_cp = ctx->tcs_vertices_per_patch;
163 uint32_t input_vertex_size = ctx->tcs_num_inputs * 16;
164 uint32_t input_patch_size = ctx->options->key.tcs.input_vertices * input_vertex_size;
165 uint32_t num_tcs_outputs = util_last_bit64(ctx->shader_info->info.tcs.outputs_written);
166 uint32_t num_tcs_patch_outputs = util_last_bit64(ctx->shader_info->info.tcs.patch_outputs_written);
167 uint32_t output_vertex_size = num_tcs_outputs * 16;
168 uint32_t pervertex_output_patch_size = ctx->tcs_vertices_per_patch * output_vertex_size;
169 uint32_t output_patch_size = pervertex_output_patch_size + num_tcs_patch_outputs * 16;
170 unsigned num_patches;
171 unsigned hardware_lds_size;
172
173 /* Ensure that we only need one wave per SIMD so we don't need to check
174 * resource usage. Also ensures that the number of tcs in and out
175 * vertices per threadgroup are at most 256.
176 */
177 num_patches = 64 / MAX2(num_tcs_input_cp, num_tcs_output_cp) * 4;
178 /* Make sure that the data fits in LDS. This assumes the shaders only
179 * use LDS for the inputs and outputs.
180 */
181 hardware_lds_size = 32768;
182
183 /* Looks like STONEY hangs if we use more than 32 KiB LDS in a single
184 * threadgroup, even though there is more than 32 KiB LDS.
185 *
186 * Test: dEQP-VK.tessellation.shader_input_output.barrier
187 */
188 if (ctx->options->chip_class >= GFX7 && ctx->options->family != CHIP_STONEY)
189 hardware_lds_size = 65536;
190
191 num_patches = MIN2(num_patches, hardware_lds_size / (input_patch_size + output_patch_size));
192 /* Make sure the output data fits in the offchip buffer */
193 num_patches = MIN2(num_patches, (ctx->options->tess_offchip_block_dw_size * 4) / output_patch_size);
194 /* Not necessary for correctness, but improves performance. The
195 * specific value is taken from the proprietary driver.
196 */
197 num_patches = MIN2(num_patches, 40);
198
199 /* GFX6 bug workaround - limit LS-HS threadgroups to only one wave. */
200 if (ctx->options->chip_class == GFX6) {
201 unsigned one_wave = ctx->options->wave_size / MAX2(num_tcs_input_cp, num_tcs_output_cp);
202 num_patches = MIN2(num_patches, one_wave);
203 }
204 return num_patches;
205 }
206
207 static unsigned
208 calculate_tess_lds_size(struct radv_shader_context *ctx)
209 {
210 unsigned num_tcs_input_cp = ctx->options->key.tcs.input_vertices;
211 unsigned num_tcs_output_cp;
212 unsigned num_tcs_outputs, num_tcs_patch_outputs;
213 unsigned input_vertex_size, output_vertex_size;
214 unsigned input_patch_size, output_patch_size;
215 unsigned pervertex_output_patch_size;
216 unsigned output_patch0_offset;
217 unsigned num_patches;
218 unsigned lds_size;
219
220 num_tcs_output_cp = ctx->tcs_vertices_per_patch;
221 num_tcs_outputs = util_last_bit64(ctx->shader_info->info.tcs.outputs_written);
222 num_tcs_patch_outputs = util_last_bit64(ctx->shader_info->info.tcs.patch_outputs_written);
223
224 input_vertex_size = ctx->tcs_num_inputs * 16;
225 output_vertex_size = num_tcs_outputs * 16;
226
227 input_patch_size = num_tcs_input_cp * input_vertex_size;
228
229 pervertex_output_patch_size = num_tcs_output_cp * output_vertex_size;
230 output_patch_size = pervertex_output_patch_size + num_tcs_patch_outputs * 16;
231
232 num_patches = ctx->tcs_num_patches;
233 output_patch0_offset = input_patch_size * num_patches;
234
235 lds_size = output_patch0_offset + output_patch_size * num_patches;
236 return lds_size;
237 }
238
239 /* Tessellation shaders pass outputs to the next shader using LDS.
240 *
241 * LS outputs = TCS inputs
242 * TCS outputs = TES inputs
243 *
244 * The LDS layout is:
245 * - TCS inputs for patch 0
246 * - TCS inputs for patch 1
247 * - TCS inputs for patch 2 = get_tcs_in_current_patch_offset (if RelPatchID==2)
248 * - ...
249 * - TCS outputs for patch 0 = get_tcs_out_patch0_offset
250 * - Per-patch TCS outputs for patch 0 = get_tcs_out_patch0_patch_data_offset
251 * - TCS outputs for patch 1
252 * - Per-patch TCS outputs for patch 1
253 * - TCS outputs for patch 2 = get_tcs_out_current_patch_offset (if RelPatchID==2)
254 * - Per-patch TCS outputs for patch 2 = get_tcs_out_current_patch_data_offset (if RelPatchID==2)
255 * - ...
256 *
257 * All three shaders VS(LS), TCS, TES share the same LDS space.
258 */
259 static LLVMValueRef
260 get_tcs_in_patch_stride(struct radv_shader_context *ctx)
261 {
262 assert (ctx->stage == MESA_SHADER_TESS_CTRL);
263 uint32_t input_vertex_size = ctx->tcs_num_inputs * 16;
264 uint32_t input_patch_size = ctx->options->key.tcs.input_vertices * input_vertex_size;
265
266 input_patch_size /= 4;
267 return LLVMConstInt(ctx->ac.i32, input_patch_size, false);
268 }
269
270 static LLVMValueRef
271 get_tcs_out_patch_stride(struct radv_shader_context *ctx)
272 {
273 uint32_t num_tcs_outputs = util_last_bit64(ctx->shader_info->info.tcs.outputs_written);
274 uint32_t num_tcs_patch_outputs = util_last_bit64(ctx->shader_info->info.tcs.patch_outputs_written);
275 uint32_t output_vertex_size = num_tcs_outputs * 16;
276 uint32_t pervertex_output_patch_size = ctx->tcs_vertices_per_patch * output_vertex_size;
277 uint32_t output_patch_size = pervertex_output_patch_size + num_tcs_patch_outputs * 16;
278 output_patch_size /= 4;
279 return LLVMConstInt(ctx->ac.i32, output_patch_size, false);
280 }
281
282 static LLVMValueRef
283 get_tcs_out_vertex_stride(struct radv_shader_context *ctx)
284 {
285 uint32_t num_tcs_outputs = util_last_bit64(ctx->shader_info->info.tcs.outputs_written);
286 uint32_t output_vertex_size = num_tcs_outputs * 16;
287 output_vertex_size /= 4;
288 return LLVMConstInt(ctx->ac.i32, output_vertex_size, false);
289 }
290
291 static LLVMValueRef
292 get_tcs_out_patch0_offset(struct radv_shader_context *ctx)
293 {
294 assert (ctx->stage == MESA_SHADER_TESS_CTRL);
295 uint32_t input_vertex_size = ctx->tcs_num_inputs * 16;
296 uint32_t input_patch_size = ctx->options->key.tcs.input_vertices * input_vertex_size;
297 uint32_t output_patch0_offset = input_patch_size;
298 unsigned num_patches = ctx->tcs_num_patches;
299
300 output_patch0_offset *= num_patches;
301 output_patch0_offset /= 4;
302 return LLVMConstInt(ctx->ac.i32, output_patch0_offset, false);
303 }
304
305 static LLVMValueRef
306 get_tcs_out_patch0_patch_data_offset(struct radv_shader_context *ctx)
307 {
308 assert (ctx->stage == MESA_SHADER_TESS_CTRL);
309 uint32_t input_vertex_size = ctx->tcs_num_inputs * 16;
310 uint32_t input_patch_size = ctx->options->key.tcs.input_vertices * input_vertex_size;
311 uint32_t output_patch0_offset = input_patch_size;
312
313 uint32_t num_tcs_outputs = util_last_bit64(ctx->shader_info->info.tcs.outputs_written);
314 uint32_t output_vertex_size = num_tcs_outputs * 16;
315 uint32_t pervertex_output_patch_size = ctx->tcs_vertices_per_patch * output_vertex_size;
316 unsigned num_patches = ctx->tcs_num_patches;
317
318 output_patch0_offset *= num_patches;
319 output_patch0_offset += pervertex_output_patch_size;
320 output_patch0_offset /= 4;
321 return LLVMConstInt(ctx->ac.i32, output_patch0_offset, false);
322 }
323
324 static LLVMValueRef
325 get_tcs_in_current_patch_offset(struct radv_shader_context *ctx)
326 {
327 LLVMValueRef patch_stride = get_tcs_in_patch_stride(ctx);
328 LLVMValueRef rel_patch_id = get_rel_patch_id(ctx);
329
330 return LLVMBuildMul(ctx->ac.builder, patch_stride, rel_patch_id, "");
331 }
332
333 static LLVMValueRef
334 get_tcs_out_current_patch_offset(struct radv_shader_context *ctx)
335 {
336 LLVMValueRef patch0_offset = get_tcs_out_patch0_offset(ctx);
337 LLVMValueRef patch_stride = get_tcs_out_patch_stride(ctx);
338 LLVMValueRef rel_patch_id = get_rel_patch_id(ctx);
339
340 return ac_build_imad(&ctx->ac, patch_stride, rel_patch_id,
341 patch0_offset);
342 }
343
344 static LLVMValueRef
345 get_tcs_out_current_patch_data_offset(struct radv_shader_context *ctx)
346 {
347 LLVMValueRef patch0_patch_data_offset =
348 get_tcs_out_patch0_patch_data_offset(ctx);
349 LLVMValueRef patch_stride = get_tcs_out_patch_stride(ctx);
350 LLVMValueRef rel_patch_id = get_rel_patch_id(ctx);
351
352 return ac_build_imad(&ctx->ac, patch_stride, rel_patch_id,
353 patch0_patch_data_offset);
354 }
355
356 #define MAX_ARGS 64
357 struct arg_info {
358 LLVMTypeRef types[MAX_ARGS];
359 LLVMValueRef *assign[MAX_ARGS];
360 uint8_t count;
361 uint8_t sgpr_count;
362 uint8_t num_sgprs_used;
363 uint8_t num_vgprs_used;
364 };
365
366 enum ac_arg_regfile {
367 ARG_SGPR,
368 ARG_VGPR,
369 };
370
371 static void
372 add_arg(struct arg_info *info, enum ac_arg_regfile regfile, LLVMTypeRef type,
373 LLVMValueRef *param_ptr)
374 {
375 assert(info->count < MAX_ARGS);
376
377 info->assign[info->count] = param_ptr;
378 info->types[info->count] = type;
379 info->count++;
380
381 if (regfile == ARG_SGPR) {
382 info->num_sgprs_used += ac_get_type_size(type) / 4;
383 info->sgpr_count++;
384 } else {
385 assert(regfile == ARG_VGPR);
386 info->num_vgprs_used += ac_get_type_size(type) / 4;
387 }
388 }
389
390 static void assign_arguments(LLVMValueRef main_function,
391 struct arg_info *info)
392 {
393 unsigned i;
394 for (i = 0; i < info->count; i++) {
395 if (info->assign[i])
396 *info->assign[i] = LLVMGetParam(main_function, i);
397 }
398 }
399
400 static LLVMValueRef
401 create_llvm_function(LLVMContextRef ctx, LLVMModuleRef module,
402 LLVMBuilderRef builder, LLVMTypeRef *return_types,
403 unsigned num_return_elems,
404 struct arg_info *args,
405 unsigned max_workgroup_size,
406 const struct radv_nir_compiler_options *options)
407 {
408 LLVMTypeRef main_function_type, ret_type;
409 LLVMBasicBlockRef main_function_body;
410
411 if (num_return_elems)
412 ret_type = LLVMStructTypeInContext(ctx, return_types,
413 num_return_elems, true);
414 else
415 ret_type = LLVMVoidTypeInContext(ctx);
416
417 /* Setup the function */
418 main_function_type =
419 LLVMFunctionType(ret_type, args->types, args->count, 0);
420 LLVMValueRef main_function =
421 LLVMAddFunction(module, "main", main_function_type);
422 main_function_body =
423 LLVMAppendBasicBlockInContext(ctx, main_function, "main_body");
424 LLVMPositionBuilderAtEnd(builder, main_function_body);
425
426 LLVMSetFunctionCallConv(main_function, RADEON_LLVM_AMDGPU_CS);
427 for (unsigned i = 0; i < args->sgpr_count; ++i) {
428 LLVMValueRef P = LLVMGetParam(main_function, i);
429
430 ac_add_function_attr(ctx, main_function, i + 1, AC_FUNC_ATTR_INREG);
431
432 if (LLVMGetTypeKind(LLVMTypeOf(P)) == LLVMPointerTypeKind) {
433 ac_add_function_attr(ctx, main_function, i + 1, AC_FUNC_ATTR_NOALIAS);
434 ac_add_attr_dereferenceable(P, UINT64_MAX);
435 }
436 }
437
438 if (options->address32_hi) {
439 ac_llvm_add_target_dep_function_attr(main_function,
440 "amdgpu-32bit-address-high-bits",
441 options->address32_hi);
442 }
443
444 ac_llvm_set_workgroup_size(main_function, max_workgroup_size);
445
446 if (options->unsafe_math) {
447 /* These were copied from some LLVM test. */
448 LLVMAddTargetDependentFunctionAttr(main_function,
449 "less-precise-fpmad",
450 "true");
451 LLVMAddTargetDependentFunctionAttr(main_function,
452 "no-infs-fp-math",
453 "true");
454 LLVMAddTargetDependentFunctionAttr(main_function,
455 "no-nans-fp-math",
456 "true");
457 LLVMAddTargetDependentFunctionAttr(main_function,
458 "unsafe-fp-math",
459 "true");
460 LLVMAddTargetDependentFunctionAttr(main_function,
461 "no-signed-zeros-fp-math",
462 "true");
463 }
464 return main_function;
465 }
466
467
468 static void
469 set_loc(struct radv_userdata_info *ud_info, uint8_t *sgpr_idx,
470 uint8_t num_sgprs)
471 {
472 ud_info->sgpr_idx = *sgpr_idx;
473 ud_info->num_sgprs = num_sgprs;
474 *sgpr_idx += num_sgprs;
475 }
476
477 static void
478 set_loc_shader(struct radv_shader_context *ctx, int idx, uint8_t *sgpr_idx,
479 uint8_t num_sgprs)
480 {
481 struct radv_userdata_info *ud_info =
482 &ctx->shader_info->user_sgprs_locs.shader_data[idx];
483 assert(ud_info);
484
485 set_loc(ud_info, sgpr_idx, num_sgprs);
486 }
487
488 static void
489 set_loc_shader_ptr(struct radv_shader_context *ctx, int idx, uint8_t *sgpr_idx)
490 {
491 bool use_32bit_pointers = idx != AC_UD_SCRATCH_RING_OFFSETS;
492
493 set_loc_shader(ctx, idx, sgpr_idx, use_32bit_pointers ? 1 : 2);
494 }
495
496 static void
497 set_loc_desc(struct radv_shader_context *ctx, int idx, uint8_t *sgpr_idx)
498 {
499 struct radv_userdata_locations *locs =
500 &ctx->shader_info->user_sgprs_locs;
501 struct radv_userdata_info *ud_info = &locs->descriptor_sets[idx];
502 assert(ud_info);
503
504 set_loc(ud_info, sgpr_idx, 1);
505
506 locs->descriptor_sets_enabled |= 1 << idx;
507 }
508
509 struct user_sgpr_info {
510 bool need_ring_offsets;
511 bool indirect_all_descriptor_sets;
512 uint8_t remaining_sgprs;
513 };
514
515 static bool needs_view_index_sgpr(struct radv_shader_context *ctx,
516 gl_shader_stage stage)
517 {
518 switch (stage) {
519 case MESA_SHADER_VERTEX:
520 if (ctx->shader_info->info.needs_multiview_view_index ||
521 (!ctx->options->key.vs_common_out.as_es && !ctx->options->key.vs_common_out.as_ls && ctx->options->key.has_multiview_view_index))
522 return true;
523 break;
524 case MESA_SHADER_TESS_EVAL:
525 if (ctx->shader_info->info.needs_multiview_view_index || (!ctx->options->key.vs_common_out.as_es && ctx->options->key.has_multiview_view_index))
526 return true;
527 break;
528 case MESA_SHADER_GEOMETRY:
529 case MESA_SHADER_TESS_CTRL:
530 if (ctx->shader_info->info.needs_multiview_view_index)
531 return true;
532 break;
533 default:
534 break;
535 }
536 return false;
537 }
538
539 static uint8_t
540 count_vs_user_sgprs(struct radv_shader_context *ctx)
541 {
542 uint8_t count = 0;
543
544 if (ctx->shader_info->info.vs.has_vertex_buffers)
545 count++;
546 count += ctx->shader_info->info.vs.needs_draw_id ? 3 : 2;
547
548 return count;
549 }
550
551 static void allocate_inline_push_consts(struct radv_shader_context *ctx,
552 struct user_sgpr_info *user_sgpr_info)
553 {
554 uint8_t remaining_sgprs = user_sgpr_info->remaining_sgprs;
555
556 /* Only supported if shaders use push constants. */
557 if (ctx->shader_info->info.min_push_constant_used == UINT8_MAX)
558 return;
559
560 /* Only supported if shaders don't have indirect push constants. */
561 if (ctx->shader_info->info.has_indirect_push_constants)
562 return;
563
564 /* Only supported for 32-bit push constants. */
565 if (!ctx->shader_info->info.has_only_32bit_push_constants)
566 return;
567
568 uint8_t num_push_consts =
569 (ctx->shader_info->info.max_push_constant_used -
570 ctx->shader_info->info.min_push_constant_used) / 4;
571
572 /* Check if the number of user SGPRs is large enough. */
573 if (num_push_consts < remaining_sgprs) {
574 ctx->shader_info->info.num_inline_push_consts = num_push_consts;
575 } else {
576 ctx->shader_info->info.num_inline_push_consts = remaining_sgprs;
577 }
578
579 /* Clamp to the maximum number of allowed inlined push constants. */
580 if (ctx->shader_info->info.num_inline_push_consts > AC_MAX_INLINE_PUSH_CONSTS)
581 ctx->shader_info->info.num_inline_push_consts = AC_MAX_INLINE_PUSH_CONSTS;
582
583 if (ctx->shader_info->info.num_inline_push_consts == num_push_consts &&
584 !ctx->shader_info->info.loads_dynamic_offsets) {
585 /* Disable the default push constants path if all constants are
586 * inlined and if shaders don't use dynamic descriptors.
587 */
588 ctx->shader_info->info.loads_push_constants = false;
589 }
590
591 ctx->shader_info->info.base_inline_push_consts =
592 ctx->shader_info->info.min_push_constant_used / 4;
593 }
594
595 static void allocate_user_sgprs(struct radv_shader_context *ctx,
596 gl_shader_stage stage,
597 bool has_previous_stage,
598 gl_shader_stage previous_stage,
599 bool needs_view_index,
600 struct user_sgpr_info *user_sgpr_info)
601 {
602 uint8_t user_sgpr_count = 0;
603
604 memset(user_sgpr_info, 0, sizeof(struct user_sgpr_info));
605
606 /* until we sort out scratch/global buffers always assign ring offsets for gs/vs/es */
607 if (stage == MESA_SHADER_GEOMETRY ||
608 stage == MESA_SHADER_VERTEX ||
609 stage == MESA_SHADER_TESS_CTRL ||
610 stage == MESA_SHADER_TESS_EVAL ||
611 ctx->is_gs_copy_shader)
612 user_sgpr_info->need_ring_offsets = true;
613
614 if (stage == MESA_SHADER_FRAGMENT &&
615 ctx->shader_info->info.ps.needs_sample_positions)
616 user_sgpr_info->need_ring_offsets = true;
617
618 /* 2 user sgprs will nearly always be allocated for scratch/rings */
619 if (ctx->options->supports_spill || user_sgpr_info->need_ring_offsets) {
620 user_sgpr_count += 2;
621 }
622
623 switch (stage) {
624 case MESA_SHADER_COMPUTE:
625 if (ctx->shader_info->info.cs.uses_grid_size)
626 user_sgpr_count += 3;
627 break;
628 case MESA_SHADER_FRAGMENT:
629 user_sgpr_count += ctx->shader_info->info.ps.needs_sample_positions;
630 break;
631 case MESA_SHADER_VERTEX:
632 if (!ctx->is_gs_copy_shader)
633 user_sgpr_count += count_vs_user_sgprs(ctx);
634 break;
635 case MESA_SHADER_TESS_CTRL:
636 if (has_previous_stage) {
637 if (previous_stage == MESA_SHADER_VERTEX)
638 user_sgpr_count += count_vs_user_sgprs(ctx);
639 }
640 break;
641 case MESA_SHADER_TESS_EVAL:
642 break;
643 case MESA_SHADER_GEOMETRY:
644 if (has_previous_stage) {
645 if (previous_stage == MESA_SHADER_VERTEX) {
646 user_sgpr_count += count_vs_user_sgprs(ctx);
647 }
648 }
649 break;
650 default:
651 break;
652 }
653
654 if (needs_view_index)
655 user_sgpr_count++;
656
657 if (ctx->shader_info->info.loads_push_constants)
658 user_sgpr_count++;
659
660 if (ctx->streamout_buffers)
661 user_sgpr_count++;
662
663 uint32_t available_sgprs = ctx->options->chip_class >= GFX9 && stage != MESA_SHADER_COMPUTE ? 32 : 16;
664 uint32_t remaining_sgprs = available_sgprs - user_sgpr_count;
665 uint32_t num_desc_set =
666 util_bitcount(ctx->shader_info->info.desc_set_used_mask);
667
668 if (remaining_sgprs < num_desc_set) {
669 user_sgpr_info->indirect_all_descriptor_sets = true;
670 user_sgpr_info->remaining_sgprs = remaining_sgprs - 1;
671 } else {
672 user_sgpr_info->remaining_sgprs = remaining_sgprs - num_desc_set;
673 }
674
675 allocate_inline_push_consts(ctx, user_sgpr_info);
676 }
677
678 static void
679 declare_global_input_sgprs(struct radv_shader_context *ctx,
680 const struct user_sgpr_info *user_sgpr_info,
681 struct arg_info *args,
682 LLVMValueRef *desc_sets)
683 {
684 LLVMTypeRef type = ac_array_in_const32_addr_space(ctx->ac.i8);
685
686 /* 1 for each descriptor set */
687 if (!user_sgpr_info->indirect_all_descriptor_sets) {
688 uint32_t mask = ctx->shader_info->info.desc_set_used_mask;
689
690 while (mask) {
691 int i = u_bit_scan(&mask);
692
693 add_arg(args, ARG_SGPR, type, &ctx->descriptor_sets[i]);
694 }
695 } else {
696 add_arg(args, ARG_SGPR, ac_array_in_const32_addr_space(type),
697 desc_sets);
698 }
699
700 if (ctx->shader_info->info.loads_push_constants) {
701 /* 1 for push constants and dynamic descriptors */
702 add_arg(args, ARG_SGPR, type, &ctx->abi.push_constants);
703 }
704
705 for (unsigned i = 0; i < ctx->shader_info->info.num_inline_push_consts; i++) {
706 add_arg(args, ARG_SGPR, ctx->ac.i32,
707 &ctx->abi.inline_push_consts[i]);
708 }
709 ctx->abi.num_inline_push_consts = ctx->shader_info->info.num_inline_push_consts;
710 ctx->abi.base_inline_push_consts = ctx->shader_info->info.base_inline_push_consts;
711
712 if (ctx->shader_info->info.so.num_outputs) {
713 add_arg(args, ARG_SGPR,
714 ac_array_in_const32_addr_space(ctx->ac.v4i32),
715 &ctx->streamout_buffers);
716 }
717 }
718
719 static void
720 declare_vs_specific_input_sgprs(struct radv_shader_context *ctx,
721 gl_shader_stage stage,
722 bool has_previous_stage,
723 gl_shader_stage previous_stage,
724 struct arg_info *args)
725 {
726 if (!ctx->is_gs_copy_shader &&
727 (stage == MESA_SHADER_VERTEX ||
728 (has_previous_stage && previous_stage == MESA_SHADER_VERTEX))) {
729 if (ctx->shader_info->info.vs.has_vertex_buffers) {
730 add_arg(args, ARG_SGPR,
731 ac_array_in_const32_addr_space(ctx->ac.v4i32),
732 &ctx->vertex_buffers);
733 }
734 add_arg(args, ARG_SGPR, ctx->ac.i32, &ctx->abi.base_vertex);
735 add_arg(args, ARG_SGPR, ctx->ac.i32, &ctx->abi.start_instance);
736 if (ctx->shader_info->info.vs.needs_draw_id) {
737 add_arg(args, ARG_SGPR, ctx->ac.i32, &ctx->abi.draw_id);
738 }
739 }
740 }
741
742 static void
743 declare_vs_input_vgprs(struct radv_shader_context *ctx, struct arg_info *args)
744 {
745 add_arg(args, ARG_VGPR, ctx->ac.i32, &ctx->abi.vertex_id);
746 if (!ctx->is_gs_copy_shader) {
747 if (ctx->options->key.vs_common_out.as_ls) {
748 add_arg(args, ARG_VGPR, ctx->ac.i32, &ctx->rel_auto_id);
749 if (ctx->ac.chip_class >= GFX10) {
750 add_arg(args, ARG_VGPR, ctx->ac.i32, NULL); /* user vgpr */
751 add_arg(args, ARG_VGPR, ctx->ac.i32, &ctx->abi.instance_id);
752 } else {
753 add_arg(args, ARG_VGPR, ctx->ac.i32, &ctx->abi.instance_id);
754 add_arg(args, ARG_VGPR, ctx->ac.i32, NULL); /* unused */
755 }
756 } else {
757 if (ctx->ac.chip_class >= GFX10) {
758 if (ctx->options->key.vs_common_out.as_ngg) {
759 add_arg(args, ARG_VGPR, ctx->ac.i32, NULL); /* user vgpr */
760 add_arg(args, ARG_VGPR, ctx->ac.i32, NULL); /* user vgpr */
761 add_arg(args, ARG_VGPR, ctx->ac.i32, &ctx->abi.instance_id);
762 } else {
763 add_arg(args, ARG_VGPR, ctx->ac.i32, NULL); /* unused */
764 add_arg(args, ARG_VGPR, ctx->ac.i32, &ctx->vs_prim_id);
765 add_arg(args, ARG_VGPR, ctx->ac.i32, &ctx->abi.instance_id);
766 }
767 } else {
768 add_arg(args, ARG_VGPR, ctx->ac.i32, &ctx->abi.instance_id);
769 add_arg(args, ARG_VGPR, ctx->ac.i32, &ctx->vs_prim_id);
770 add_arg(args, ARG_VGPR, ctx->ac.i32, NULL); /* unused */
771 }
772 }
773 }
774 }
775
776 static void
777 declare_streamout_sgprs(struct radv_shader_context *ctx, gl_shader_stage stage,
778 struct arg_info *args)
779 {
780 int i;
781
782 /* Streamout SGPRs. */
783 if (ctx->shader_info->info.so.num_outputs) {
784 assert(stage == MESA_SHADER_VERTEX ||
785 stage == MESA_SHADER_TESS_EVAL);
786
787 if (stage != MESA_SHADER_TESS_EVAL) {
788 add_arg(args, ARG_SGPR, ctx->ac.i32, &ctx->streamout_config);
789 } else {
790 args->assign[args->count - 1] = &ctx->streamout_config;
791 args->types[args->count - 1] = ctx->ac.i32;
792 }
793
794 add_arg(args, ARG_SGPR, ctx->ac.i32, &ctx->streamout_write_idx);
795 }
796
797 /* A streamout buffer offset is loaded if the stride is non-zero. */
798 for (i = 0; i < 4; i++) {
799 if (!ctx->shader_info->info.so.strides[i])
800 continue;
801
802 add_arg(args, ARG_SGPR, ctx->ac.i32, &ctx->streamout_offset[i]);
803 }
804 }
805
806 static void
807 declare_tes_input_vgprs(struct radv_shader_context *ctx, struct arg_info *args)
808 {
809 add_arg(args, ARG_VGPR, ctx->ac.f32, &ctx->tes_u);
810 add_arg(args, ARG_VGPR, ctx->ac.f32, &ctx->tes_v);
811 add_arg(args, ARG_VGPR, ctx->ac.i32, &ctx->tes_rel_patch_id);
812 add_arg(args, ARG_VGPR, ctx->ac.i32, &ctx->abi.tes_patch_id);
813 }
814
815 static void
816 set_global_input_locs(struct radv_shader_context *ctx,
817 const struct user_sgpr_info *user_sgpr_info,
818 LLVMValueRef desc_sets, uint8_t *user_sgpr_idx)
819 {
820 uint32_t mask = ctx->shader_info->info.desc_set_used_mask;
821
822 if (!user_sgpr_info->indirect_all_descriptor_sets) {
823 while (mask) {
824 int i = u_bit_scan(&mask);
825
826 set_loc_desc(ctx, i, user_sgpr_idx);
827 }
828 } else {
829 set_loc_shader_ptr(ctx, AC_UD_INDIRECT_DESCRIPTOR_SETS,
830 user_sgpr_idx);
831
832 while (mask) {
833 int i = u_bit_scan(&mask);
834
835 ctx->descriptor_sets[i] =
836 ac_build_load_to_sgpr(&ctx->ac, desc_sets,
837 LLVMConstInt(ctx->ac.i32, i, false));
838
839 }
840
841 ctx->shader_info->need_indirect_descriptor_sets = true;
842 }
843
844 if (ctx->shader_info->info.loads_push_constants) {
845 set_loc_shader_ptr(ctx, AC_UD_PUSH_CONSTANTS, user_sgpr_idx);
846 }
847
848 if (ctx->shader_info->info.num_inline_push_consts) {
849 set_loc_shader(ctx, AC_UD_INLINE_PUSH_CONSTANTS, user_sgpr_idx,
850 ctx->shader_info->info.num_inline_push_consts);
851 }
852
853 if (ctx->streamout_buffers) {
854 set_loc_shader_ptr(ctx, AC_UD_STREAMOUT_BUFFERS,
855 user_sgpr_idx);
856 }
857 }
858
859 static void
860 set_vs_specific_input_locs(struct radv_shader_context *ctx,
861 gl_shader_stage stage, bool has_previous_stage,
862 gl_shader_stage previous_stage,
863 uint8_t *user_sgpr_idx)
864 {
865 if (!ctx->is_gs_copy_shader &&
866 (stage == MESA_SHADER_VERTEX ||
867 (has_previous_stage && previous_stage == MESA_SHADER_VERTEX))) {
868 if (ctx->shader_info->info.vs.has_vertex_buffers) {
869 set_loc_shader_ptr(ctx, AC_UD_VS_VERTEX_BUFFERS,
870 user_sgpr_idx);
871 }
872
873 unsigned vs_num = 2;
874 if (ctx->shader_info->info.vs.needs_draw_id)
875 vs_num++;
876
877 set_loc_shader(ctx, AC_UD_VS_BASE_VERTEX_START_INSTANCE,
878 user_sgpr_idx, vs_num);
879 }
880 }
881
882 static void set_llvm_calling_convention(LLVMValueRef func,
883 gl_shader_stage stage)
884 {
885 enum radeon_llvm_calling_convention calling_conv;
886
887 switch (stage) {
888 case MESA_SHADER_VERTEX:
889 case MESA_SHADER_TESS_EVAL:
890 calling_conv = RADEON_LLVM_AMDGPU_VS;
891 break;
892 case MESA_SHADER_GEOMETRY:
893 calling_conv = RADEON_LLVM_AMDGPU_GS;
894 break;
895 case MESA_SHADER_TESS_CTRL:
896 calling_conv = RADEON_LLVM_AMDGPU_HS;
897 break;
898 case MESA_SHADER_FRAGMENT:
899 calling_conv = RADEON_LLVM_AMDGPU_PS;
900 break;
901 case MESA_SHADER_COMPUTE:
902 calling_conv = RADEON_LLVM_AMDGPU_CS;
903 break;
904 default:
905 unreachable("Unhandle shader type");
906 }
907
908 LLVMSetFunctionCallConv(func, calling_conv);
909 }
910
911 /* Returns whether the stage is a stage that can be directly before the GS */
912 static bool is_pre_gs_stage(gl_shader_stage stage)
913 {
914 return stage == MESA_SHADER_VERTEX || stage == MESA_SHADER_TESS_EVAL;
915 }
916
917 static void create_function(struct radv_shader_context *ctx,
918 gl_shader_stage stage,
919 bool has_previous_stage,
920 gl_shader_stage previous_stage)
921 {
922 uint8_t user_sgpr_idx;
923 struct user_sgpr_info user_sgpr_info;
924 struct arg_info args = {};
925 LLVMValueRef desc_sets;
926 bool needs_view_index = needs_view_index_sgpr(ctx, stage);
927
928 if (ctx->ac.chip_class >= GFX10) {
929 if (is_pre_gs_stage(stage) && ctx->options->key.vs_common_out.as_ngg) {
930 /* On GFX10, VS is merged into GS for NGG. */
931 previous_stage = stage;
932 stage = MESA_SHADER_GEOMETRY;
933 has_previous_stage = true;
934 }
935 }
936
937 allocate_user_sgprs(ctx, stage, has_previous_stage,
938 previous_stage, needs_view_index, &user_sgpr_info);
939
940 if (user_sgpr_info.need_ring_offsets && !ctx->options->supports_spill) {
941 add_arg(&args, ARG_SGPR, ac_array_in_const_addr_space(ctx->ac.v4i32),
942 &ctx->ring_offsets);
943 }
944
945 switch (stage) {
946 case MESA_SHADER_COMPUTE:
947 declare_global_input_sgprs(ctx, &user_sgpr_info, &args,
948 &desc_sets);
949
950 if (ctx->shader_info->info.cs.uses_grid_size) {
951 add_arg(&args, ARG_SGPR, ctx->ac.v3i32,
952 &ctx->abi.num_work_groups);
953 }
954
955 for (int i = 0; i < 3; i++) {
956 ctx->abi.workgroup_ids[i] = NULL;
957 if (ctx->shader_info->info.cs.uses_block_id[i]) {
958 add_arg(&args, ARG_SGPR, ctx->ac.i32,
959 &ctx->abi.workgroup_ids[i]);
960 }
961 }
962
963 if (ctx->shader_info->info.cs.uses_local_invocation_idx)
964 add_arg(&args, ARG_SGPR, ctx->ac.i32, &ctx->abi.tg_size);
965 add_arg(&args, ARG_VGPR, ctx->ac.v3i32,
966 &ctx->abi.local_invocation_ids);
967 break;
968 case MESA_SHADER_VERTEX:
969 declare_global_input_sgprs(ctx, &user_sgpr_info, &args,
970 &desc_sets);
971
972 declare_vs_specific_input_sgprs(ctx, stage, has_previous_stage,
973 previous_stage, &args);
974
975 if (needs_view_index)
976 add_arg(&args, ARG_SGPR, ctx->ac.i32,
977 &ctx->abi.view_index);
978 if (ctx->options->key.vs_common_out.as_es) {
979 add_arg(&args, ARG_SGPR, ctx->ac.i32,
980 &ctx->es2gs_offset);
981 } else if (ctx->options->key.vs_common_out.as_ls) {
982 /* no extra parameters */
983 } else {
984 declare_streamout_sgprs(ctx, stage, &args);
985 }
986
987 declare_vs_input_vgprs(ctx, &args);
988 break;
989 case MESA_SHADER_TESS_CTRL:
990 if (has_previous_stage) {
991 // First 6 system regs
992 add_arg(&args, ARG_SGPR, ctx->ac.i32, &ctx->oc_lds);
993 add_arg(&args, ARG_SGPR, ctx->ac.i32,
994 &ctx->merged_wave_info);
995 add_arg(&args, ARG_SGPR, ctx->ac.i32,
996 &ctx->tess_factor_offset);
997
998 add_arg(&args, ARG_SGPR, ctx->ac.i32, NULL); // scratch offset
999 add_arg(&args, ARG_SGPR, ctx->ac.i32, NULL); // unknown
1000 add_arg(&args, ARG_SGPR, ctx->ac.i32, NULL); // unknown
1001
1002 declare_global_input_sgprs(ctx, &user_sgpr_info, &args,
1003 &desc_sets);
1004
1005 declare_vs_specific_input_sgprs(ctx, stage,
1006 has_previous_stage,
1007 previous_stage, &args);
1008
1009 if (needs_view_index)
1010 add_arg(&args, ARG_SGPR, ctx->ac.i32,
1011 &ctx->abi.view_index);
1012
1013 add_arg(&args, ARG_VGPR, ctx->ac.i32,
1014 &ctx->abi.tcs_patch_id);
1015 add_arg(&args, ARG_VGPR, ctx->ac.i32,
1016 &ctx->abi.tcs_rel_ids);
1017
1018 declare_vs_input_vgprs(ctx, &args);
1019 } else {
1020 declare_global_input_sgprs(ctx, &user_sgpr_info, &args,
1021 &desc_sets);
1022
1023 if (needs_view_index)
1024 add_arg(&args, ARG_SGPR, ctx->ac.i32,
1025 &ctx->abi.view_index);
1026
1027 add_arg(&args, ARG_SGPR, ctx->ac.i32, &ctx->oc_lds);
1028 add_arg(&args, ARG_SGPR, ctx->ac.i32,
1029 &ctx->tess_factor_offset);
1030 add_arg(&args, ARG_VGPR, ctx->ac.i32,
1031 &ctx->abi.tcs_patch_id);
1032 add_arg(&args, ARG_VGPR, ctx->ac.i32,
1033 &ctx->abi.tcs_rel_ids);
1034 }
1035 break;
1036 case MESA_SHADER_TESS_EVAL:
1037 declare_global_input_sgprs(ctx, &user_sgpr_info, &args,
1038 &desc_sets);
1039
1040 if (needs_view_index)
1041 add_arg(&args, ARG_SGPR, ctx->ac.i32,
1042 &ctx->abi.view_index);
1043
1044 if (ctx->options->key.vs_common_out.as_es) {
1045 add_arg(&args, ARG_SGPR, ctx->ac.i32, &ctx->oc_lds);
1046 add_arg(&args, ARG_SGPR, ctx->ac.i32, NULL);
1047 add_arg(&args, ARG_SGPR, ctx->ac.i32,
1048 &ctx->es2gs_offset);
1049 } else {
1050 add_arg(&args, ARG_SGPR, ctx->ac.i32, NULL);
1051 declare_streamout_sgprs(ctx, stage, &args);
1052 add_arg(&args, ARG_SGPR, ctx->ac.i32, &ctx->oc_lds);
1053 }
1054 declare_tes_input_vgprs(ctx, &args);
1055 break;
1056 case MESA_SHADER_GEOMETRY:
1057 if (has_previous_stage) {
1058 // First 6 system regs
1059 if (ctx->options->key.vs_common_out.as_ngg) {
1060 add_arg(&args, ARG_SGPR, ctx->ac.i32,
1061 &ctx->gs_tg_info);
1062 } else {
1063 add_arg(&args, ARG_SGPR, ctx->ac.i32,
1064 &ctx->gs2vs_offset);
1065 }
1066
1067 add_arg(&args, ARG_SGPR, ctx->ac.i32,
1068 &ctx->merged_wave_info);
1069 add_arg(&args, ARG_SGPR, ctx->ac.i32, &ctx->oc_lds);
1070
1071 add_arg(&args, ARG_SGPR, ctx->ac.i32, NULL); // scratch offset
1072 add_arg(&args, ARG_SGPR, ctx->ac.i32, NULL); // unknown
1073 add_arg(&args, ARG_SGPR, ctx->ac.i32, NULL); // unknown
1074
1075 declare_global_input_sgprs(ctx, &user_sgpr_info, &args,
1076 &desc_sets);
1077
1078 if (previous_stage != MESA_SHADER_TESS_EVAL) {
1079 declare_vs_specific_input_sgprs(ctx, stage,
1080 has_previous_stage,
1081 previous_stage,
1082 &args);
1083 }
1084
1085 if (needs_view_index)
1086 add_arg(&args, ARG_SGPR, ctx->ac.i32,
1087 &ctx->abi.view_index);
1088
1089 add_arg(&args, ARG_VGPR, ctx->ac.i32,
1090 &ctx->gs_vtx_offset[0]);
1091 add_arg(&args, ARG_VGPR, ctx->ac.i32,
1092 &ctx->gs_vtx_offset[2]);
1093 add_arg(&args, ARG_VGPR, ctx->ac.i32,
1094 &ctx->abi.gs_prim_id);
1095 add_arg(&args, ARG_VGPR, ctx->ac.i32,
1096 &ctx->abi.gs_invocation_id);
1097 add_arg(&args, ARG_VGPR, ctx->ac.i32,
1098 &ctx->gs_vtx_offset[4]);
1099
1100 if (previous_stage == MESA_SHADER_VERTEX) {
1101 declare_vs_input_vgprs(ctx, &args);
1102 } else {
1103 declare_tes_input_vgprs(ctx, &args);
1104 }
1105 } else {
1106 declare_global_input_sgprs(ctx, &user_sgpr_info, &args,
1107 &desc_sets);
1108
1109 if (needs_view_index)
1110 add_arg(&args, ARG_SGPR, ctx->ac.i32,
1111 &ctx->abi.view_index);
1112
1113 add_arg(&args, ARG_SGPR, ctx->ac.i32, &ctx->gs2vs_offset);
1114 add_arg(&args, ARG_SGPR, ctx->ac.i32, &ctx->gs_wave_id);
1115 add_arg(&args, ARG_VGPR, ctx->ac.i32,
1116 &ctx->gs_vtx_offset[0]);
1117 add_arg(&args, ARG_VGPR, ctx->ac.i32,
1118 &ctx->gs_vtx_offset[1]);
1119 add_arg(&args, ARG_VGPR, ctx->ac.i32,
1120 &ctx->abi.gs_prim_id);
1121 add_arg(&args, ARG_VGPR, ctx->ac.i32,
1122 &ctx->gs_vtx_offset[2]);
1123 add_arg(&args, ARG_VGPR, ctx->ac.i32,
1124 &ctx->gs_vtx_offset[3]);
1125 add_arg(&args, ARG_VGPR, ctx->ac.i32,
1126 &ctx->gs_vtx_offset[4]);
1127 add_arg(&args, ARG_VGPR, ctx->ac.i32,
1128 &ctx->gs_vtx_offset[5]);
1129 add_arg(&args, ARG_VGPR, ctx->ac.i32,
1130 &ctx->abi.gs_invocation_id);
1131 }
1132 break;
1133 case MESA_SHADER_FRAGMENT:
1134 declare_global_input_sgprs(ctx, &user_sgpr_info, &args,
1135 &desc_sets);
1136
1137 add_arg(&args, ARG_SGPR, ctx->ac.i32, &ctx->abi.prim_mask);
1138 add_arg(&args, ARG_VGPR, ctx->ac.v2i32, &ctx->abi.persp_sample);
1139 add_arg(&args, ARG_VGPR, ctx->ac.v2i32, &ctx->abi.persp_center);
1140 add_arg(&args, ARG_VGPR, ctx->ac.v2i32, &ctx->abi.persp_centroid);
1141 add_arg(&args, ARG_VGPR, ctx->ac.v3i32, NULL); /* persp pull model */
1142 add_arg(&args, ARG_VGPR, ctx->ac.v2i32, &ctx->abi.linear_sample);
1143 add_arg(&args, ARG_VGPR, ctx->ac.v2i32, &ctx->abi.linear_center);
1144 add_arg(&args, ARG_VGPR, ctx->ac.v2i32, &ctx->abi.linear_centroid);
1145 add_arg(&args, ARG_VGPR, ctx->ac.f32, NULL); /* line stipple tex */
1146 add_arg(&args, ARG_VGPR, ctx->ac.f32, &ctx->abi.frag_pos[0]);
1147 add_arg(&args, ARG_VGPR, ctx->ac.f32, &ctx->abi.frag_pos[1]);
1148 add_arg(&args, ARG_VGPR, ctx->ac.f32, &ctx->abi.frag_pos[2]);
1149 add_arg(&args, ARG_VGPR, ctx->ac.f32, &ctx->abi.frag_pos[3]);
1150 add_arg(&args, ARG_VGPR, ctx->ac.i32, &ctx->abi.front_face);
1151 add_arg(&args, ARG_VGPR, ctx->ac.i32, &ctx->abi.ancillary);
1152 add_arg(&args, ARG_VGPR, ctx->ac.i32, &ctx->abi.sample_coverage);
1153 add_arg(&args, ARG_VGPR, ctx->ac.i32, NULL); /* fixed pt */
1154 break;
1155 default:
1156 unreachable("Shader stage not implemented");
1157 }
1158
1159 ctx->main_function = create_llvm_function(
1160 ctx->context, ctx->ac.module, ctx->ac.builder, NULL, 0, &args,
1161 ctx->max_workgroup_size, ctx->options);
1162 set_llvm_calling_convention(ctx->main_function, stage);
1163
1164
1165 ctx->shader_info->num_input_vgprs = 0;
1166 ctx->shader_info->num_input_sgprs = ctx->options->supports_spill ? 2 : 0;
1167
1168 ctx->shader_info->num_input_sgprs += args.num_sgprs_used;
1169
1170 if (ctx->stage != MESA_SHADER_FRAGMENT)
1171 ctx->shader_info->num_input_vgprs = args.num_vgprs_used;
1172
1173 assign_arguments(ctx->main_function, &args);
1174
1175 user_sgpr_idx = 0;
1176
1177 if (ctx->options->supports_spill || user_sgpr_info.need_ring_offsets) {
1178 set_loc_shader_ptr(ctx, AC_UD_SCRATCH_RING_OFFSETS,
1179 &user_sgpr_idx);
1180 if (ctx->options->supports_spill) {
1181 ctx->ring_offsets = ac_build_intrinsic(&ctx->ac, "llvm.amdgcn.implicit.buffer.ptr",
1182 LLVMPointerType(ctx->ac.i8, AC_ADDR_SPACE_CONST),
1183 NULL, 0, AC_FUNC_ATTR_READNONE);
1184 ctx->ring_offsets = LLVMBuildBitCast(ctx->ac.builder, ctx->ring_offsets,
1185 ac_array_in_const_addr_space(ctx->ac.v4i32), "");
1186 }
1187 }
1188
1189 /* For merged shaders the user SGPRs start at 8, with 8 system SGPRs in front (including
1190 * the rw_buffers at s0/s1. With user SGPR0 = s8, lets restart the count from 0 */
1191 if (has_previous_stage)
1192 user_sgpr_idx = 0;
1193
1194 set_global_input_locs(ctx, &user_sgpr_info, desc_sets, &user_sgpr_idx);
1195
1196 switch (stage) {
1197 case MESA_SHADER_COMPUTE:
1198 if (ctx->shader_info->info.cs.uses_grid_size) {
1199 set_loc_shader(ctx, AC_UD_CS_GRID_SIZE,
1200 &user_sgpr_idx, 3);
1201 }
1202 break;
1203 case MESA_SHADER_VERTEX:
1204 set_vs_specific_input_locs(ctx, stage, has_previous_stage,
1205 previous_stage, &user_sgpr_idx);
1206 if (ctx->abi.view_index)
1207 set_loc_shader(ctx, AC_UD_VIEW_INDEX, &user_sgpr_idx, 1);
1208 break;
1209 case MESA_SHADER_TESS_CTRL:
1210 set_vs_specific_input_locs(ctx, stage, has_previous_stage,
1211 previous_stage, &user_sgpr_idx);
1212 if (ctx->abi.view_index)
1213 set_loc_shader(ctx, AC_UD_VIEW_INDEX, &user_sgpr_idx, 1);
1214 break;
1215 case MESA_SHADER_TESS_EVAL:
1216 if (ctx->abi.view_index)
1217 set_loc_shader(ctx, AC_UD_VIEW_INDEX, &user_sgpr_idx, 1);
1218 break;
1219 case MESA_SHADER_GEOMETRY:
1220 if (has_previous_stage) {
1221 if (previous_stage == MESA_SHADER_VERTEX)
1222 set_vs_specific_input_locs(ctx, stage,
1223 has_previous_stage,
1224 previous_stage,
1225 &user_sgpr_idx);
1226 }
1227 if (ctx->abi.view_index)
1228 set_loc_shader(ctx, AC_UD_VIEW_INDEX, &user_sgpr_idx, 1);
1229 break;
1230 case MESA_SHADER_FRAGMENT:
1231 break;
1232 default:
1233 unreachable("Shader stage not implemented");
1234 }
1235
1236 if (stage == MESA_SHADER_TESS_CTRL ||
1237 (stage == MESA_SHADER_VERTEX && ctx->options->key.vs_common_out.as_ls) ||
1238 /* GFX9 has the ESGS ring buffer in LDS. */
1239 (stage == MESA_SHADER_GEOMETRY && has_previous_stage)) {
1240 ac_declare_lds_as_pointer(&ctx->ac);
1241 }
1242
1243 ctx->shader_info->num_user_sgprs = user_sgpr_idx;
1244 }
1245
1246
1247 static LLVMValueRef
1248 radv_load_resource(struct ac_shader_abi *abi, LLVMValueRef index,
1249 unsigned desc_set, unsigned binding)
1250 {
1251 struct radv_shader_context *ctx = radv_shader_context_from_abi(abi);
1252 LLVMValueRef desc_ptr = ctx->descriptor_sets[desc_set];
1253 struct radv_pipeline_layout *pipeline_layout = ctx->options->layout;
1254 struct radv_descriptor_set_layout *layout = pipeline_layout->set[desc_set].layout;
1255 unsigned base_offset = layout->binding[binding].offset;
1256 LLVMValueRef offset, stride;
1257
1258 if (layout->binding[binding].type == VK_DESCRIPTOR_TYPE_UNIFORM_BUFFER_DYNAMIC ||
1259 layout->binding[binding].type == VK_DESCRIPTOR_TYPE_STORAGE_BUFFER_DYNAMIC) {
1260 unsigned idx = pipeline_layout->set[desc_set].dynamic_offset_start +
1261 layout->binding[binding].dynamic_offset_offset;
1262 desc_ptr = ctx->abi.push_constants;
1263 base_offset = pipeline_layout->push_constant_size + 16 * idx;
1264 stride = LLVMConstInt(ctx->ac.i32, 16, false);
1265 } else
1266 stride = LLVMConstInt(ctx->ac.i32, layout->binding[binding].size, false);
1267
1268 offset = LLVMConstInt(ctx->ac.i32, base_offset, false);
1269
1270 if (layout->binding[binding].type != VK_DESCRIPTOR_TYPE_INLINE_UNIFORM_BLOCK_EXT) {
1271 offset = ac_build_imad(&ctx->ac, index, stride, offset);
1272 }
1273
1274 desc_ptr = LLVMBuildGEP(ctx->ac.builder, desc_ptr, &offset, 1, "");
1275 desc_ptr = ac_cast_ptr(&ctx->ac, desc_ptr, ctx->ac.v4i32);
1276 LLVMSetMetadata(desc_ptr, ctx->ac.uniform_md_kind, ctx->ac.empty_md);
1277
1278 if (layout->binding[binding].type == VK_DESCRIPTOR_TYPE_INLINE_UNIFORM_BLOCK_EXT) {
1279 uint32_t desc_type = S_008F0C_DST_SEL_X(V_008F0C_SQ_SEL_X) |
1280 S_008F0C_DST_SEL_Y(V_008F0C_SQ_SEL_Y) |
1281 S_008F0C_DST_SEL_Z(V_008F0C_SQ_SEL_Z) |
1282 S_008F0C_DST_SEL_W(V_008F0C_SQ_SEL_W);
1283
1284 if (ctx->ac.chip_class >= GFX10) {
1285 desc_type |= S_008F0C_FORMAT(V_008F0C_IMG_FORMAT_32_FLOAT) |
1286 S_008F0C_OOB_SELECT(3) |
1287 S_008F0C_RESOURCE_LEVEL(1);
1288 } else {
1289 desc_type |= S_008F0C_NUM_FORMAT(V_008F0C_BUF_NUM_FORMAT_FLOAT) |
1290 S_008F0C_DATA_FORMAT(V_008F0C_BUF_DATA_FORMAT_32);
1291 }
1292
1293 LLVMValueRef desc_components[4] = {
1294 LLVMBuildPtrToInt(ctx->ac.builder, desc_ptr, ctx->ac.intptr, ""),
1295 LLVMConstInt(ctx->ac.i32, S_008F04_BASE_ADDRESS_HI(ctx->options->address32_hi), false),
1296 /* High limit to support variable sizes. */
1297 LLVMConstInt(ctx->ac.i32, 0xffffffff, false),
1298 LLVMConstInt(ctx->ac.i32, desc_type, false),
1299 };
1300
1301 return ac_build_gather_values(&ctx->ac, desc_components, 4);
1302 }
1303
1304 return desc_ptr;
1305 }
1306
1307
1308 /* The offchip buffer layout for TCS->TES is
1309 *
1310 * - attribute 0 of patch 0 vertex 0
1311 * - attribute 0 of patch 0 vertex 1
1312 * - attribute 0 of patch 0 vertex 2
1313 * ...
1314 * - attribute 0 of patch 1 vertex 0
1315 * - attribute 0 of patch 1 vertex 1
1316 * ...
1317 * - attribute 1 of patch 0 vertex 0
1318 * - attribute 1 of patch 0 vertex 1
1319 * ...
1320 * - per patch attribute 0 of patch 0
1321 * - per patch attribute 0 of patch 1
1322 * ...
1323 *
1324 * Note that every attribute has 4 components.
1325 */
1326 static LLVMValueRef get_non_vertex_index_offset(struct radv_shader_context *ctx)
1327 {
1328 uint32_t num_patches = ctx->tcs_num_patches;
1329 uint32_t num_tcs_outputs;
1330 if (ctx->stage == MESA_SHADER_TESS_CTRL)
1331 num_tcs_outputs = util_last_bit64(ctx->shader_info->info.tcs.outputs_written);
1332 else
1333 num_tcs_outputs = ctx->options->key.tes.tcs_num_outputs;
1334
1335 uint32_t output_vertex_size = num_tcs_outputs * 16;
1336 uint32_t pervertex_output_patch_size = ctx->tcs_vertices_per_patch * output_vertex_size;
1337
1338 return LLVMConstInt(ctx->ac.i32, pervertex_output_patch_size * num_patches, false);
1339 }
1340
1341 static LLVMValueRef calc_param_stride(struct radv_shader_context *ctx,
1342 LLVMValueRef vertex_index)
1343 {
1344 LLVMValueRef param_stride;
1345 if (vertex_index)
1346 param_stride = LLVMConstInt(ctx->ac.i32, ctx->tcs_vertices_per_patch * ctx->tcs_num_patches, false);
1347 else
1348 param_stride = LLVMConstInt(ctx->ac.i32, ctx->tcs_num_patches, false);
1349 return param_stride;
1350 }
1351
1352 static LLVMValueRef get_tcs_tes_buffer_address(struct radv_shader_context *ctx,
1353 LLVMValueRef vertex_index,
1354 LLVMValueRef param_index)
1355 {
1356 LLVMValueRef base_addr;
1357 LLVMValueRef param_stride, constant16;
1358 LLVMValueRef rel_patch_id = get_rel_patch_id(ctx);
1359 LLVMValueRef vertices_per_patch = LLVMConstInt(ctx->ac.i32, ctx->tcs_vertices_per_patch, false);
1360 constant16 = LLVMConstInt(ctx->ac.i32, 16, false);
1361 param_stride = calc_param_stride(ctx, vertex_index);
1362 if (vertex_index) {
1363 base_addr = ac_build_imad(&ctx->ac, rel_patch_id,
1364 vertices_per_patch, vertex_index);
1365 } else {
1366 base_addr = rel_patch_id;
1367 }
1368
1369 base_addr = LLVMBuildAdd(ctx->ac.builder, base_addr,
1370 LLVMBuildMul(ctx->ac.builder, param_index,
1371 param_stride, ""), "");
1372
1373 base_addr = LLVMBuildMul(ctx->ac.builder, base_addr, constant16, "");
1374
1375 if (!vertex_index) {
1376 LLVMValueRef patch_data_offset = get_non_vertex_index_offset(ctx);
1377
1378 base_addr = LLVMBuildAdd(ctx->ac.builder, base_addr,
1379 patch_data_offset, "");
1380 }
1381 return base_addr;
1382 }
1383
1384 static LLVMValueRef get_tcs_tes_buffer_address_params(struct radv_shader_context *ctx,
1385 unsigned param,
1386 unsigned const_index,
1387 bool is_compact,
1388 LLVMValueRef vertex_index,
1389 LLVMValueRef indir_index)
1390 {
1391 LLVMValueRef param_index;
1392
1393 if (indir_index)
1394 param_index = LLVMBuildAdd(ctx->ac.builder, LLVMConstInt(ctx->ac.i32, param, false),
1395 indir_index, "");
1396 else {
1397 if (const_index && !is_compact)
1398 param += const_index;
1399 param_index = LLVMConstInt(ctx->ac.i32, param, false);
1400 }
1401 return get_tcs_tes_buffer_address(ctx, vertex_index, param_index);
1402 }
1403
1404 static LLVMValueRef
1405 get_dw_address(struct radv_shader_context *ctx,
1406 LLVMValueRef dw_addr,
1407 unsigned param,
1408 unsigned const_index,
1409 bool compact_const_index,
1410 LLVMValueRef vertex_index,
1411 LLVMValueRef stride,
1412 LLVMValueRef indir_index)
1413
1414 {
1415
1416 if (vertex_index) {
1417 dw_addr = LLVMBuildAdd(ctx->ac.builder, dw_addr,
1418 LLVMBuildMul(ctx->ac.builder,
1419 vertex_index,
1420 stride, ""), "");
1421 }
1422
1423 if (indir_index)
1424 dw_addr = LLVMBuildAdd(ctx->ac.builder, dw_addr,
1425 LLVMBuildMul(ctx->ac.builder, indir_index,
1426 LLVMConstInt(ctx->ac.i32, 4, false), ""), "");
1427 else if (const_index && !compact_const_index)
1428 dw_addr = LLVMBuildAdd(ctx->ac.builder, dw_addr,
1429 LLVMConstInt(ctx->ac.i32, const_index * 4, false), "");
1430
1431 dw_addr = LLVMBuildAdd(ctx->ac.builder, dw_addr,
1432 LLVMConstInt(ctx->ac.i32, param * 4, false), "");
1433
1434 if (const_index && compact_const_index)
1435 dw_addr = LLVMBuildAdd(ctx->ac.builder, dw_addr,
1436 LLVMConstInt(ctx->ac.i32, const_index, false), "");
1437 return dw_addr;
1438 }
1439
1440 static LLVMValueRef
1441 load_tcs_varyings(struct ac_shader_abi *abi,
1442 LLVMTypeRef type,
1443 LLVMValueRef vertex_index,
1444 LLVMValueRef indir_index,
1445 unsigned const_index,
1446 unsigned location,
1447 unsigned driver_location,
1448 unsigned component,
1449 unsigned num_components,
1450 bool is_patch,
1451 bool is_compact,
1452 bool load_input)
1453 {
1454 struct radv_shader_context *ctx = radv_shader_context_from_abi(abi);
1455 LLVMValueRef dw_addr, stride;
1456 LLVMValueRef value[4], result;
1457 unsigned param = shader_io_get_unique_index(location);
1458
1459 if (load_input) {
1460 uint32_t input_vertex_size = (ctx->tcs_num_inputs * 16) / 4;
1461 stride = LLVMConstInt(ctx->ac.i32, input_vertex_size, false);
1462 dw_addr = get_tcs_in_current_patch_offset(ctx);
1463 } else {
1464 if (!is_patch) {
1465 stride = get_tcs_out_vertex_stride(ctx);
1466 dw_addr = get_tcs_out_current_patch_offset(ctx);
1467 } else {
1468 dw_addr = get_tcs_out_current_patch_data_offset(ctx);
1469 stride = NULL;
1470 }
1471 }
1472
1473 dw_addr = get_dw_address(ctx, dw_addr, param, const_index, is_compact, vertex_index, stride,
1474 indir_index);
1475
1476 for (unsigned i = 0; i < num_components + component; i++) {
1477 value[i] = ac_lds_load(&ctx->ac, dw_addr);
1478 dw_addr = LLVMBuildAdd(ctx->ac.builder, dw_addr,
1479 ctx->ac.i32_1, "");
1480 }
1481 result = ac_build_varying_gather_values(&ctx->ac, value, num_components, component);
1482 return result;
1483 }
1484
1485 static void
1486 store_tcs_output(struct ac_shader_abi *abi,
1487 const nir_variable *var,
1488 LLVMValueRef vertex_index,
1489 LLVMValueRef param_index,
1490 unsigned const_index,
1491 LLVMValueRef src,
1492 unsigned writemask)
1493 {
1494 struct radv_shader_context *ctx = radv_shader_context_from_abi(abi);
1495 const unsigned location = var->data.location;
1496 unsigned component = var->data.location_frac;
1497 const bool is_patch = var->data.patch;
1498 const bool is_compact = var->data.compact;
1499 LLVMValueRef dw_addr;
1500 LLVMValueRef stride = NULL;
1501 LLVMValueRef buf_addr = NULL;
1502 unsigned param;
1503 bool store_lds = true;
1504
1505 if (is_patch) {
1506 if (!(ctx->tcs_patch_outputs_read & (1U << (location - VARYING_SLOT_PATCH0))))
1507 store_lds = false;
1508 } else {
1509 if (!(ctx->tcs_outputs_read & (1ULL << location)))
1510 store_lds = false;
1511 }
1512
1513 param = shader_io_get_unique_index(location);
1514 if ((location == VARYING_SLOT_CLIP_DIST0 || location == VARYING_SLOT_CLIP_DIST1) && is_compact) {
1515 const_index += component;
1516 component = 0;
1517
1518 if (const_index >= 4) {
1519 const_index -= 4;
1520 param++;
1521 }
1522 }
1523
1524 if (!is_patch) {
1525 stride = get_tcs_out_vertex_stride(ctx);
1526 dw_addr = get_tcs_out_current_patch_offset(ctx);
1527 } else {
1528 dw_addr = get_tcs_out_current_patch_data_offset(ctx);
1529 }
1530
1531 dw_addr = get_dw_address(ctx, dw_addr, param, const_index, is_compact, vertex_index, stride,
1532 param_index);
1533 buf_addr = get_tcs_tes_buffer_address_params(ctx, param, const_index, is_compact,
1534 vertex_index, param_index);
1535
1536 bool is_tess_factor = false;
1537 if (location == VARYING_SLOT_TESS_LEVEL_INNER ||
1538 location == VARYING_SLOT_TESS_LEVEL_OUTER)
1539 is_tess_factor = true;
1540
1541 unsigned base = is_compact ? const_index : 0;
1542 for (unsigned chan = 0; chan < 8; chan++) {
1543 if (!(writemask & (1 << chan)))
1544 continue;
1545 LLVMValueRef value = ac_llvm_extract_elem(&ctx->ac, src, chan - component);
1546 value = ac_to_integer(&ctx->ac, value);
1547 value = LLVMBuildZExtOrBitCast(ctx->ac.builder, value, ctx->ac.i32, "");
1548
1549 if (store_lds || is_tess_factor) {
1550 LLVMValueRef dw_addr_chan =
1551 LLVMBuildAdd(ctx->ac.builder, dw_addr,
1552 LLVMConstInt(ctx->ac.i32, chan, false), "");
1553 ac_lds_store(&ctx->ac, dw_addr_chan, value);
1554 }
1555
1556 if (!is_tess_factor && writemask != 0xF)
1557 ac_build_buffer_store_dword(&ctx->ac, ctx->hs_ring_tess_offchip, value, 1,
1558 buf_addr, ctx->oc_lds,
1559 4 * (base + chan), ac_glc, false);
1560 }
1561
1562 if (writemask == 0xF) {
1563 ac_build_buffer_store_dword(&ctx->ac, ctx->hs_ring_tess_offchip, src, 4,
1564 buf_addr, ctx->oc_lds,
1565 (base * 4), ac_glc, false);
1566 }
1567 }
1568
1569 static LLVMValueRef
1570 load_tes_input(struct ac_shader_abi *abi,
1571 LLVMTypeRef type,
1572 LLVMValueRef vertex_index,
1573 LLVMValueRef param_index,
1574 unsigned const_index,
1575 unsigned location,
1576 unsigned driver_location,
1577 unsigned component,
1578 unsigned num_components,
1579 bool is_patch,
1580 bool is_compact,
1581 bool load_input)
1582 {
1583 struct radv_shader_context *ctx = radv_shader_context_from_abi(abi);
1584 LLVMValueRef buf_addr;
1585 LLVMValueRef result;
1586 unsigned param = shader_io_get_unique_index(location);
1587
1588 if ((location == VARYING_SLOT_CLIP_DIST0 || location == VARYING_SLOT_CLIP_DIST1) && is_compact) {
1589 const_index += component;
1590 component = 0;
1591 if (const_index >= 4) {
1592 const_index -= 4;
1593 param++;
1594 }
1595 }
1596
1597 buf_addr = get_tcs_tes_buffer_address_params(ctx, param, const_index,
1598 is_compact, vertex_index, param_index);
1599
1600 LLVMValueRef comp_offset = LLVMConstInt(ctx->ac.i32, component * 4, false);
1601 buf_addr = LLVMBuildAdd(ctx->ac.builder, buf_addr, comp_offset, "");
1602
1603 result = ac_build_buffer_load(&ctx->ac, ctx->hs_ring_tess_offchip, num_components, NULL,
1604 buf_addr, ctx->oc_lds, is_compact ? (4 * const_index) : 0, ac_glc, true, false);
1605 result = ac_trim_vector(&ctx->ac, result, num_components);
1606 return result;
1607 }
1608
1609 static LLVMValueRef
1610 load_gs_input(struct ac_shader_abi *abi,
1611 unsigned location,
1612 unsigned driver_location,
1613 unsigned component,
1614 unsigned num_components,
1615 unsigned vertex_index,
1616 unsigned const_index,
1617 LLVMTypeRef type)
1618 {
1619 struct radv_shader_context *ctx = radv_shader_context_from_abi(abi);
1620 LLVMValueRef vtx_offset;
1621 unsigned param, vtx_offset_param;
1622 LLVMValueRef value[4], result;
1623
1624 vtx_offset_param = vertex_index;
1625 assert(vtx_offset_param < 6);
1626 vtx_offset = LLVMBuildMul(ctx->ac.builder, ctx->gs_vtx_offset[vtx_offset_param],
1627 LLVMConstInt(ctx->ac.i32, 4, false), "");
1628
1629 param = shader_io_get_unique_index(location);
1630
1631 for (unsigned i = component; i < num_components + component; i++) {
1632 if (ctx->ac.chip_class >= GFX9) {
1633 LLVMValueRef dw_addr = ctx->gs_vtx_offset[vtx_offset_param];
1634 dw_addr = LLVMBuildAdd(ctx->ac.builder, dw_addr,
1635 LLVMConstInt(ctx->ac.i32, param * 4 + i + const_index, 0), "");
1636 value[i] = ac_lds_load(&ctx->ac, dw_addr);
1637 } else {
1638 LLVMValueRef soffset =
1639 LLVMConstInt(ctx->ac.i32,
1640 (param * 4 + i + const_index) * 256,
1641 false);
1642
1643 value[i] = ac_build_buffer_load(&ctx->ac,
1644 ctx->esgs_ring, 1,
1645 ctx->ac.i32_0,
1646 vtx_offset, soffset,
1647 0, ac_glc, true, false);
1648 }
1649
1650 if (ac_get_type_size(type) == 2) {
1651 value[i] = LLVMBuildBitCast(ctx->ac.builder, value[i], ctx->ac.i32, "");
1652 value[i] = LLVMBuildTrunc(ctx->ac.builder, value[i], ctx->ac.i16, "");
1653 }
1654 value[i] = LLVMBuildBitCast(ctx->ac.builder, value[i], type, "");
1655 }
1656 result = ac_build_varying_gather_values(&ctx->ac, value, num_components, component);
1657 result = ac_to_integer(&ctx->ac, result);
1658 return result;
1659 }
1660
1661
1662 static void radv_emit_kill(struct ac_shader_abi *abi, LLVMValueRef visible)
1663 {
1664 struct radv_shader_context *ctx = radv_shader_context_from_abi(abi);
1665 ac_build_kill_if_false(&ctx->ac, visible);
1666 }
1667
1668 static uint32_t
1669 radv_get_sample_pos_offset(uint32_t num_samples)
1670 {
1671 uint32_t sample_pos_offset = 0;
1672
1673 switch (num_samples) {
1674 case 2:
1675 sample_pos_offset = 1;
1676 break;
1677 case 4:
1678 sample_pos_offset = 3;
1679 break;
1680 case 8:
1681 sample_pos_offset = 7;
1682 break;
1683 default:
1684 break;
1685 }
1686 return sample_pos_offset;
1687 }
1688
1689 static LLVMValueRef load_sample_position(struct ac_shader_abi *abi,
1690 LLVMValueRef sample_id)
1691 {
1692 struct radv_shader_context *ctx = radv_shader_context_from_abi(abi);
1693
1694 LLVMValueRef result;
1695 LLVMValueRef index = LLVMConstInt(ctx->ac.i32, RING_PS_SAMPLE_POSITIONS, false);
1696 LLVMValueRef ptr = LLVMBuildGEP(ctx->ac.builder, ctx->ring_offsets, &index, 1, "");
1697
1698 ptr = LLVMBuildBitCast(ctx->ac.builder, ptr,
1699 ac_array_in_const_addr_space(ctx->ac.v2f32), "");
1700
1701 uint32_t sample_pos_offset =
1702 radv_get_sample_pos_offset(ctx->options->key.fs.num_samples);
1703
1704 sample_id =
1705 LLVMBuildAdd(ctx->ac.builder, sample_id,
1706 LLVMConstInt(ctx->ac.i32, sample_pos_offset, false), "");
1707 result = ac_build_load_invariant(&ctx->ac, ptr, sample_id);
1708
1709 return result;
1710 }
1711
1712
1713 static LLVMValueRef load_sample_mask_in(struct ac_shader_abi *abi)
1714 {
1715 struct radv_shader_context *ctx = radv_shader_context_from_abi(abi);
1716 uint8_t log2_ps_iter_samples;
1717
1718 if (ctx->shader_info->info.ps.force_persample) {
1719 log2_ps_iter_samples =
1720 util_logbase2(ctx->options->key.fs.num_samples);
1721 } else {
1722 log2_ps_iter_samples = ctx->options->key.fs.log2_ps_iter_samples;
1723 }
1724
1725 /* The bit pattern matches that used by fixed function fragment
1726 * processing. */
1727 static const uint16_t ps_iter_masks[] = {
1728 0xffff, /* not used */
1729 0x5555,
1730 0x1111,
1731 0x0101,
1732 0x0001,
1733 };
1734 assert(log2_ps_iter_samples < ARRAY_SIZE(ps_iter_masks));
1735
1736 uint32_t ps_iter_mask = ps_iter_masks[log2_ps_iter_samples];
1737
1738 LLVMValueRef result, sample_id;
1739 sample_id = ac_unpack_param(&ctx->ac, abi->ancillary, 8, 4);
1740 sample_id = LLVMBuildShl(ctx->ac.builder, LLVMConstInt(ctx->ac.i32, ps_iter_mask, false), sample_id, "");
1741 result = LLVMBuildAnd(ctx->ac.builder, sample_id, abi->sample_coverage, "");
1742 return result;
1743 }
1744
1745
1746 static void gfx10_ngg_gs_emit_vertex(struct radv_shader_context *ctx,
1747 unsigned stream,
1748 LLVMValueRef *addrs);
1749
1750 static void
1751 visit_emit_vertex(struct ac_shader_abi *abi, unsigned stream, LLVMValueRef *addrs)
1752 {
1753 LLVMValueRef gs_next_vertex;
1754 LLVMValueRef can_emit;
1755 unsigned offset = 0;
1756 struct radv_shader_context *ctx = radv_shader_context_from_abi(abi);
1757
1758 if (ctx->options->key.vs_common_out.as_ngg) {
1759 gfx10_ngg_gs_emit_vertex(ctx, stream, addrs);
1760 return;
1761 }
1762
1763 /* Write vertex attribute values to GSVS ring */
1764 gs_next_vertex = LLVMBuildLoad(ctx->ac.builder,
1765 ctx->gs_next_vertex[stream],
1766 "");
1767
1768 /* If this thread has already emitted the declared maximum number of
1769 * vertices, kill it: excessive vertex emissions are not supposed to
1770 * have any effect, and GS threads have no externally observable
1771 * effects other than emitting vertices.
1772 */
1773 can_emit = LLVMBuildICmp(ctx->ac.builder, LLVMIntULT, gs_next_vertex,
1774 LLVMConstInt(ctx->ac.i32, ctx->gs_max_out_vertices, false), "");
1775 ac_build_kill_if_false(&ctx->ac, can_emit);
1776
1777 for (unsigned i = 0; i < AC_LLVM_MAX_OUTPUTS; ++i) {
1778 unsigned output_usage_mask =
1779 ctx->shader_info->info.gs.output_usage_mask[i];
1780 uint8_t output_stream =
1781 ctx->shader_info->info.gs.output_streams[i];
1782 LLVMValueRef *out_ptr = &addrs[i * 4];
1783 int length = util_last_bit(output_usage_mask);
1784
1785 if (!(ctx->output_mask & (1ull << i)) ||
1786 output_stream != stream)
1787 continue;
1788
1789 for (unsigned j = 0; j < length; j++) {
1790 if (!(output_usage_mask & (1 << j)))
1791 continue;
1792
1793 LLVMValueRef out_val = LLVMBuildLoad(ctx->ac.builder,
1794 out_ptr[j], "");
1795 LLVMValueRef voffset =
1796 LLVMConstInt(ctx->ac.i32, offset *
1797 ctx->gs_max_out_vertices, false);
1798
1799 offset++;
1800
1801 voffset = LLVMBuildAdd(ctx->ac.builder, voffset, gs_next_vertex, "");
1802 voffset = LLVMBuildMul(ctx->ac.builder, voffset, LLVMConstInt(ctx->ac.i32, 4, false), "");
1803
1804 out_val = ac_to_integer(&ctx->ac, out_val);
1805 out_val = LLVMBuildZExtOrBitCast(ctx->ac.builder, out_val, ctx->ac.i32, "");
1806
1807 ac_build_buffer_store_dword(&ctx->ac,
1808 ctx->gsvs_ring[stream],
1809 out_val, 1,
1810 voffset, ctx->gs2vs_offset, 0,
1811 ac_glc | ac_slc, true);
1812 }
1813 }
1814
1815 gs_next_vertex = LLVMBuildAdd(ctx->ac.builder, gs_next_vertex,
1816 ctx->ac.i32_1, "");
1817 LLVMBuildStore(ctx->ac.builder, gs_next_vertex, ctx->gs_next_vertex[stream]);
1818
1819 ac_build_sendmsg(&ctx->ac,
1820 AC_SENDMSG_GS_OP_EMIT | AC_SENDMSG_GS | (stream << 8),
1821 ctx->gs_wave_id);
1822 }
1823
1824 static void
1825 visit_end_primitive(struct ac_shader_abi *abi, unsigned stream)
1826 {
1827 struct radv_shader_context *ctx = radv_shader_context_from_abi(abi);
1828
1829 if (ctx->options->key.vs_common_out.as_ngg) {
1830 LLVMBuildStore(ctx->ac.builder, ctx->ac.i32_0, ctx->gs_curprim_verts[stream]);
1831 return;
1832 }
1833
1834 ac_build_sendmsg(&ctx->ac, AC_SENDMSG_GS_OP_CUT | AC_SENDMSG_GS | (stream << 8), ctx->gs_wave_id);
1835 }
1836
1837 static LLVMValueRef
1838 load_tess_coord(struct ac_shader_abi *abi)
1839 {
1840 struct radv_shader_context *ctx = radv_shader_context_from_abi(abi);
1841
1842 LLVMValueRef coord[4] = {
1843 ctx->tes_u,
1844 ctx->tes_v,
1845 ctx->ac.f32_0,
1846 ctx->ac.f32_0,
1847 };
1848
1849 if (ctx->tes_primitive_mode == GL_TRIANGLES)
1850 coord[2] = LLVMBuildFSub(ctx->ac.builder, ctx->ac.f32_1,
1851 LLVMBuildFAdd(ctx->ac.builder, coord[0], coord[1], ""), "");
1852
1853 return ac_build_gather_values(&ctx->ac, coord, 3);
1854 }
1855
1856 static LLVMValueRef
1857 load_patch_vertices_in(struct ac_shader_abi *abi)
1858 {
1859 struct radv_shader_context *ctx = radv_shader_context_from_abi(abi);
1860 return LLVMConstInt(ctx->ac.i32, ctx->options->key.tcs.input_vertices, false);
1861 }
1862
1863
1864 static LLVMValueRef radv_load_base_vertex(struct ac_shader_abi *abi)
1865 {
1866 return abi->base_vertex;
1867 }
1868
1869 static LLVMValueRef radv_load_ssbo(struct ac_shader_abi *abi,
1870 LLVMValueRef buffer_ptr, bool write)
1871 {
1872 struct radv_shader_context *ctx = radv_shader_context_from_abi(abi);
1873 LLVMValueRef result;
1874
1875 LLVMSetMetadata(buffer_ptr, ctx->ac.uniform_md_kind, ctx->ac.empty_md);
1876
1877 result = LLVMBuildLoad(ctx->ac.builder, buffer_ptr, "");
1878 LLVMSetMetadata(result, ctx->ac.invariant_load_md_kind, ctx->ac.empty_md);
1879
1880 return result;
1881 }
1882
1883 static LLVMValueRef radv_load_ubo(struct ac_shader_abi *abi, LLVMValueRef buffer_ptr)
1884 {
1885 struct radv_shader_context *ctx = radv_shader_context_from_abi(abi);
1886 LLVMValueRef result;
1887
1888 if (LLVMGetTypeKind(LLVMTypeOf(buffer_ptr)) != LLVMPointerTypeKind) {
1889 /* Do not load the descriptor for inlined uniform blocks. */
1890 return buffer_ptr;
1891 }
1892
1893 LLVMSetMetadata(buffer_ptr, ctx->ac.uniform_md_kind, ctx->ac.empty_md);
1894
1895 result = LLVMBuildLoad(ctx->ac.builder, buffer_ptr, "");
1896 LLVMSetMetadata(result, ctx->ac.invariant_load_md_kind, ctx->ac.empty_md);
1897
1898 return result;
1899 }
1900
1901 static LLVMValueRef radv_get_sampler_desc(struct ac_shader_abi *abi,
1902 unsigned descriptor_set,
1903 unsigned base_index,
1904 unsigned constant_index,
1905 LLVMValueRef index,
1906 enum ac_descriptor_type desc_type,
1907 bool image, bool write,
1908 bool bindless)
1909 {
1910 struct radv_shader_context *ctx = radv_shader_context_from_abi(abi);
1911 LLVMValueRef list = ctx->descriptor_sets[descriptor_set];
1912 struct radv_descriptor_set_layout *layout = ctx->options->layout->set[descriptor_set].layout;
1913 struct radv_descriptor_set_binding_layout *binding = layout->binding + base_index;
1914 unsigned offset = binding->offset;
1915 unsigned stride = binding->size;
1916 unsigned type_size;
1917 LLVMBuilderRef builder = ctx->ac.builder;
1918 LLVMTypeRef type;
1919
1920 assert(base_index < layout->binding_count);
1921
1922 switch (desc_type) {
1923 case AC_DESC_IMAGE:
1924 type = ctx->ac.v8i32;
1925 type_size = 32;
1926 break;
1927 case AC_DESC_FMASK:
1928 type = ctx->ac.v8i32;
1929 offset += 32;
1930 type_size = 32;
1931 break;
1932 case AC_DESC_SAMPLER:
1933 type = ctx->ac.v4i32;
1934 if (binding->type == VK_DESCRIPTOR_TYPE_COMBINED_IMAGE_SAMPLER) {
1935 offset += radv_combined_image_descriptor_sampler_offset(binding);
1936 }
1937
1938 type_size = 16;
1939 break;
1940 case AC_DESC_BUFFER:
1941 type = ctx->ac.v4i32;
1942 type_size = 16;
1943 break;
1944 case AC_DESC_PLANE_0:
1945 case AC_DESC_PLANE_1:
1946 case AC_DESC_PLANE_2:
1947 type = ctx->ac.v8i32;
1948 type_size = 32;
1949 offset += 32 * (desc_type - AC_DESC_PLANE_0);
1950 break;
1951 default:
1952 unreachable("invalid desc_type\n");
1953 }
1954
1955 offset += constant_index * stride;
1956
1957 if (desc_type == AC_DESC_SAMPLER && binding->immutable_samplers_offset &&
1958 (!index || binding->immutable_samplers_equal)) {
1959 if (binding->immutable_samplers_equal)
1960 constant_index = 0;
1961
1962 const uint32_t *samplers = radv_immutable_samplers(layout, binding);
1963
1964 LLVMValueRef constants[] = {
1965 LLVMConstInt(ctx->ac.i32, samplers[constant_index * 4 + 0], 0),
1966 LLVMConstInt(ctx->ac.i32, samplers[constant_index * 4 + 1], 0),
1967 LLVMConstInt(ctx->ac.i32, samplers[constant_index * 4 + 2], 0),
1968 LLVMConstInt(ctx->ac.i32, samplers[constant_index * 4 + 3], 0),
1969 };
1970 return ac_build_gather_values(&ctx->ac, constants, 4);
1971 }
1972
1973 assert(stride % type_size == 0);
1974
1975 LLVMValueRef adjusted_index = index;
1976 if (!adjusted_index)
1977 adjusted_index = ctx->ac.i32_0;
1978
1979 adjusted_index = LLVMBuildMul(builder, adjusted_index, LLVMConstInt(ctx->ac.i32, stride / type_size, 0), "");
1980
1981 LLVMValueRef val_offset = LLVMConstInt(ctx->ac.i32, offset, 0);
1982 list = LLVMBuildGEP(builder, list, &val_offset, 1, "");
1983 list = LLVMBuildPointerCast(builder, list,
1984 ac_array_in_const32_addr_space(type), "");
1985
1986 LLVMValueRef descriptor = ac_build_load_to_sgpr(&ctx->ac, list, adjusted_index);
1987
1988 /* 3 plane formats always have same size and format for plane 1 & 2, so
1989 * use the tail from plane 1 so that we can store only the first 16 bytes
1990 * of the last plane. */
1991 if (desc_type == AC_DESC_PLANE_2) {
1992 LLVMValueRef descriptor2 = radv_get_sampler_desc(abi, descriptor_set, base_index, constant_index, index, AC_DESC_PLANE_1,image, write, bindless);
1993
1994 LLVMValueRef components[8];
1995 for (unsigned i = 0; i < 4; ++i)
1996 components[i] = ac_llvm_extract_elem(&ctx->ac, descriptor, i);
1997
1998 for (unsigned i = 4; i < 8; ++i)
1999 components[i] = ac_llvm_extract_elem(&ctx->ac, descriptor2, i);
2000 descriptor = ac_build_gather_values(&ctx->ac, components, 8);
2001 }
2002
2003 return descriptor;
2004 }
2005
2006 /* For 2_10_10_10 formats the alpha is handled as unsigned by pre-vega HW.
2007 * so we may need to fix it up. */
2008 static LLVMValueRef
2009 adjust_vertex_fetch_alpha(struct radv_shader_context *ctx,
2010 unsigned adjustment,
2011 LLVMValueRef alpha)
2012 {
2013 if (adjustment == RADV_ALPHA_ADJUST_NONE)
2014 return alpha;
2015
2016 LLVMValueRef c30 = LLVMConstInt(ctx->ac.i32, 30, 0);
2017
2018 alpha = LLVMBuildBitCast(ctx->ac.builder, alpha, ctx->ac.f32, "");
2019
2020 if (adjustment == RADV_ALPHA_ADJUST_SSCALED)
2021 alpha = LLVMBuildFPToUI(ctx->ac.builder, alpha, ctx->ac.i32, "");
2022 else
2023 alpha = ac_to_integer(&ctx->ac, alpha);
2024
2025 /* For the integer-like cases, do a natural sign extension.
2026 *
2027 * For the SNORM case, the values are 0.0, 0.333, 0.666, 1.0
2028 * and happen to contain 0, 1, 2, 3 as the two LSBs of the
2029 * exponent.
2030 */
2031 alpha = LLVMBuildShl(ctx->ac.builder, alpha,
2032 adjustment == RADV_ALPHA_ADJUST_SNORM ?
2033 LLVMConstInt(ctx->ac.i32, 7, 0) : c30, "");
2034 alpha = LLVMBuildAShr(ctx->ac.builder, alpha, c30, "");
2035
2036 /* Convert back to the right type. */
2037 if (adjustment == RADV_ALPHA_ADJUST_SNORM) {
2038 LLVMValueRef clamp;
2039 LLVMValueRef neg_one = LLVMConstReal(ctx->ac.f32, -1.0);
2040 alpha = LLVMBuildSIToFP(ctx->ac.builder, alpha, ctx->ac.f32, "");
2041 clamp = LLVMBuildFCmp(ctx->ac.builder, LLVMRealULT, alpha, neg_one, "");
2042 alpha = LLVMBuildSelect(ctx->ac.builder, clamp, neg_one, alpha, "");
2043 } else if (adjustment == RADV_ALPHA_ADJUST_SSCALED) {
2044 alpha = LLVMBuildSIToFP(ctx->ac.builder, alpha, ctx->ac.f32, "");
2045 }
2046
2047 return LLVMBuildBitCast(ctx->ac.builder, alpha, ctx->ac.i32, "");
2048 }
2049
2050 static unsigned
2051 get_num_channels_from_data_format(unsigned data_format)
2052 {
2053 switch (data_format) {
2054 case V_008F0C_BUF_DATA_FORMAT_8:
2055 case V_008F0C_BUF_DATA_FORMAT_16:
2056 case V_008F0C_BUF_DATA_FORMAT_32:
2057 return 1;
2058 case V_008F0C_BUF_DATA_FORMAT_8_8:
2059 case V_008F0C_BUF_DATA_FORMAT_16_16:
2060 case V_008F0C_BUF_DATA_FORMAT_32_32:
2061 return 2;
2062 case V_008F0C_BUF_DATA_FORMAT_10_11_11:
2063 case V_008F0C_BUF_DATA_FORMAT_11_11_10:
2064 case V_008F0C_BUF_DATA_FORMAT_32_32_32:
2065 return 3;
2066 case V_008F0C_BUF_DATA_FORMAT_8_8_8_8:
2067 case V_008F0C_BUF_DATA_FORMAT_10_10_10_2:
2068 case V_008F0C_BUF_DATA_FORMAT_2_10_10_10:
2069 case V_008F0C_BUF_DATA_FORMAT_16_16_16_16:
2070 case V_008F0C_BUF_DATA_FORMAT_32_32_32_32:
2071 return 4;
2072 default:
2073 break;
2074 }
2075
2076 return 4;
2077 }
2078
2079 static LLVMValueRef
2080 radv_fixup_vertex_input_fetches(struct radv_shader_context *ctx,
2081 LLVMValueRef value,
2082 unsigned num_channels,
2083 bool is_float)
2084 {
2085 LLVMValueRef zero = is_float ? ctx->ac.f32_0 : ctx->ac.i32_0;
2086 LLVMValueRef one = is_float ? ctx->ac.f32_1 : ctx->ac.i32_1;
2087 LLVMValueRef chan[4];
2088
2089 if (LLVMGetTypeKind(LLVMTypeOf(value)) == LLVMVectorTypeKind) {
2090 unsigned vec_size = LLVMGetVectorSize(LLVMTypeOf(value));
2091
2092 if (num_channels == 4 && num_channels == vec_size)
2093 return value;
2094
2095 num_channels = MIN2(num_channels, vec_size);
2096
2097 for (unsigned i = 0; i < num_channels; i++)
2098 chan[i] = ac_llvm_extract_elem(&ctx->ac, value, i);
2099 } else {
2100 if (num_channels) {
2101 assert(num_channels == 1);
2102 chan[0] = value;
2103 }
2104 }
2105
2106 for (unsigned i = num_channels; i < 4; i++) {
2107 chan[i] = i == 3 ? one : zero;
2108 chan[i] = ac_to_integer(&ctx->ac, chan[i]);
2109 }
2110
2111 return ac_build_gather_values(&ctx->ac, chan, 4);
2112 }
2113
2114 static void
2115 handle_vs_input_decl(struct radv_shader_context *ctx,
2116 struct nir_variable *variable)
2117 {
2118 LLVMValueRef t_list_ptr = ctx->vertex_buffers;
2119 LLVMValueRef t_offset;
2120 LLVMValueRef t_list;
2121 LLVMValueRef input;
2122 LLVMValueRef buffer_index;
2123 unsigned attrib_count = glsl_count_attribute_slots(variable->type, true);
2124 uint8_t input_usage_mask =
2125 ctx->shader_info->info.vs.input_usage_mask[variable->data.location];
2126 unsigned num_input_channels = util_last_bit(input_usage_mask);
2127
2128 variable->data.driver_location = variable->data.location * 4;
2129
2130 enum glsl_base_type type = glsl_get_base_type(variable->type);
2131 for (unsigned i = 0; i < attrib_count; ++i) {
2132 LLVMValueRef output[4];
2133 unsigned attrib_index = variable->data.location + i - VERT_ATTRIB_GENERIC0;
2134 unsigned attrib_format = ctx->options->key.vs.vertex_attribute_formats[attrib_index];
2135 unsigned data_format = attrib_format & 0x0f;
2136 unsigned num_format = (attrib_format >> 4) & 0x07;
2137 bool is_float = num_format != V_008F0C_BUF_NUM_FORMAT_UINT &&
2138 num_format != V_008F0C_BUF_NUM_FORMAT_SINT;
2139
2140 if (ctx->options->key.vs.instance_rate_inputs & (1u << attrib_index)) {
2141 uint32_t divisor = ctx->options->key.vs.instance_rate_divisors[attrib_index];
2142
2143 if (divisor) {
2144 buffer_index = ctx->abi.instance_id;
2145
2146 if (divisor != 1) {
2147 buffer_index = LLVMBuildUDiv(ctx->ac.builder, buffer_index,
2148 LLVMConstInt(ctx->ac.i32, divisor, 0), "");
2149 }
2150 } else {
2151 buffer_index = ctx->ac.i32_0;
2152 }
2153
2154 buffer_index = LLVMBuildAdd(ctx->ac.builder, ctx->abi.start_instance, buffer_index, "");
2155 } else
2156 buffer_index = LLVMBuildAdd(ctx->ac.builder, ctx->abi.vertex_id,
2157 ctx->abi.base_vertex, "");
2158
2159 /* Adjust the number of channels to load based on the vertex
2160 * attribute format.
2161 */
2162 unsigned num_format_channels = get_num_channels_from_data_format(data_format);
2163 unsigned num_channels = MIN2(num_input_channels, num_format_channels);
2164 unsigned attrib_binding = ctx->options->key.vs.vertex_attribute_bindings[attrib_index];
2165 unsigned attrib_offset = ctx->options->key.vs.vertex_attribute_offsets[attrib_index];
2166 unsigned attrib_stride = ctx->options->key.vs.vertex_attribute_strides[attrib_index];
2167
2168 if (ctx->options->key.vs.post_shuffle & (1 << attrib_index)) {
2169 /* Always load, at least, 3 channels for formats that
2170 * need to be shuffled because X<->Z.
2171 */
2172 num_channels = MAX2(num_channels, 3);
2173 }
2174
2175 if (attrib_stride != 0 && attrib_offset > attrib_stride) {
2176 LLVMValueRef buffer_offset =
2177 LLVMConstInt(ctx->ac.i32,
2178 attrib_offset / attrib_stride, false);
2179
2180 buffer_index = LLVMBuildAdd(ctx->ac.builder,
2181 buffer_index,
2182 buffer_offset, "");
2183
2184 attrib_offset = attrib_offset % attrib_stride;
2185 }
2186
2187 t_offset = LLVMConstInt(ctx->ac.i32, attrib_binding, false);
2188 t_list = ac_build_load_to_sgpr(&ctx->ac, t_list_ptr, t_offset);
2189
2190 input = ac_build_struct_tbuffer_load(&ctx->ac, t_list,
2191 buffer_index,
2192 LLVMConstInt(ctx->ac.i32, attrib_offset, false),
2193 ctx->ac.i32_0, ctx->ac.i32_0,
2194 num_channels,
2195 data_format, num_format, 0, true);
2196
2197 if (ctx->options->key.vs.post_shuffle & (1 << attrib_index)) {
2198 LLVMValueRef c[4];
2199 c[0] = ac_llvm_extract_elem(&ctx->ac, input, 2);
2200 c[1] = ac_llvm_extract_elem(&ctx->ac, input, 1);
2201 c[2] = ac_llvm_extract_elem(&ctx->ac, input, 0);
2202 c[3] = ac_llvm_extract_elem(&ctx->ac, input, 3);
2203
2204 input = ac_build_gather_values(&ctx->ac, c, 4);
2205 }
2206
2207 input = radv_fixup_vertex_input_fetches(ctx, input, num_channels,
2208 is_float);
2209
2210 for (unsigned chan = 0; chan < 4; chan++) {
2211 LLVMValueRef llvm_chan = LLVMConstInt(ctx->ac.i32, chan, false);
2212 output[chan] = LLVMBuildExtractElement(ctx->ac.builder, input, llvm_chan, "");
2213 if (type == GLSL_TYPE_FLOAT16) {
2214 output[chan] = LLVMBuildBitCast(ctx->ac.builder, output[chan], ctx->ac.f32, "");
2215 output[chan] = LLVMBuildFPTrunc(ctx->ac.builder, output[chan], ctx->ac.f16, "");
2216 }
2217 }
2218
2219 unsigned alpha_adjust = (ctx->options->key.vs.alpha_adjust >> (attrib_index * 2)) & 3;
2220 output[3] = adjust_vertex_fetch_alpha(ctx, alpha_adjust, output[3]);
2221
2222 for (unsigned chan = 0; chan < 4; chan++) {
2223 output[chan] = ac_to_integer(&ctx->ac, output[chan]);
2224 if (type == GLSL_TYPE_UINT16 || type == GLSL_TYPE_INT16)
2225 output[chan] = LLVMBuildTrunc(ctx->ac.builder, output[chan], ctx->ac.i16, "");
2226
2227 ctx->inputs[ac_llvm_reg_index_soa(variable->data.location + i, chan)] = output[chan];
2228 }
2229 }
2230 }
2231
2232 static void
2233 handle_vs_inputs(struct radv_shader_context *ctx,
2234 struct nir_shader *nir) {
2235 nir_foreach_variable(variable, &nir->inputs)
2236 handle_vs_input_decl(ctx, variable);
2237 }
2238
2239 static void
2240 prepare_interp_optimize(struct radv_shader_context *ctx,
2241 struct nir_shader *nir)
2242 {
2243 bool uses_center = false;
2244 bool uses_centroid = false;
2245 nir_foreach_variable(variable, &nir->inputs) {
2246 if (glsl_get_base_type(glsl_without_array(variable->type)) != GLSL_TYPE_FLOAT ||
2247 variable->data.sample)
2248 continue;
2249
2250 if (variable->data.centroid)
2251 uses_centroid = true;
2252 else
2253 uses_center = true;
2254 }
2255
2256 if (uses_center && uses_centroid) {
2257 LLVMValueRef sel = LLVMBuildICmp(ctx->ac.builder, LLVMIntSLT, ctx->abi.prim_mask, ctx->ac.i32_0, "");
2258 ctx->abi.persp_centroid = LLVMBuildSelect(ctx->ac.builder, sel, ctx->abi.persp_center, ctx->abi.persp_centroid, "");
2259 ctx->abi.linear_centroid = LLVMBuildSelect(ctx->ac.builder, sel, ctx->abi.linear_center, ctx->abi.linear_centroid, "");
2260 }
2261 }
2262
2263 static void
2264 scan_shader_output_decl(struct radv_shader_context *ctx,
2265 struct nir_variable *variable,
2266 struct nir_shader *shader,
2267 gl_shader_stage stage)
2268 {
2269 int idx = variable->data.location + variable->data.index;
2270 unsigned attrib_count = glsl_count_attribute_slots(variable->type, false);
2271 uint64_t mask_attribs;
2272
2273 variable->data.driver_location = idx * 4;
2274
2275 /* tess ctrl has it's own load/store paths for outputs */
2276 if (stage == MESA_SHADER_TESS_CTRL)
2277 return;
2278
2279 if (variable->data.compact) {
2280 unsigned component_count = variable->data.location_frac +
2281 glsl_get_length(variable->type);
2282 attrib_count = (component_count + 3) / 4;
2283 }
2284
2285 mask_attribs = ((1ull << attrib_count) - 1) << idx;
2286 if (stage == MESA_SHADER_VERTEX ||
2287 stage == MESA_SHADER_TESS_EVAL ||
2288 stage == MESA_SHADER_GEOMETRY) {
2289 if (idx == VARYING_SLOT_CLIP_DIST0) {
2290 if (stage == MESA_SHADER_VERTEX) {
2291 ctx->shader_info->vs.outinfo.clip_dist_mask = (1 << shader->info.clip_distance_array_size) - 1;
2292 ctx->shader_info->vs.outinfo.cull_dist_mask = (1 << shader->info.cull_distance_array_size) - 1;
2293 ctx->shader_info->vs.outinfo.cull_dist_mask <<= shader->info.clip_distance_array_size;
2294 }
2295 if (stage == MESA_SHADER_TESS_EVAL) {
2296 ctx->shader_info->tes.outinfo.clip_dist_mask = (1 << shader->info.clip_distance_array_size) - 1;
2297 ctx->shader_info->tes.outinfo.cull_dist_mask = (1 << shader->info.cull_distance_array_size) - 1;
2298 ctx->shader_info->tes.outinfo.cull_dist_mask <<= shader->info.clip_distance_array_size;
2299 }
2300 if (stage == MESA_SHADER_GEOMETRY) {
2301 ctx->shader_info->vs.outinfo.clip_dist_mask = (1 << shader->info.clip_distance_array_size) - 1;
2302 ctx->shader_info->vs.outinfo.cull_dist_mask = (1 << shader->info.cull_distance_array_size) - 1;
2303 ctx->shader_info->vs.outinfo.cull_dist_mask <<= shader->info.clip_distance_array_size;
2304 }
2305 }
2306 }
2307
2308 ctx->output_mask |= mask_attribs;
2309 }
2310
2311
2312 /* Initialize arguments for the shader export intrinsic */
2313 static void
2314 si_llvm_init_export_args(struct radv_shader_context *ctx,
2315 LLVMValueRef *values,
2316 unsigned enabled_channels,
2317 unsigned target,
2318 struct ac_export_args *args)
2319 {
2320 /* Specify the channels that are enabled. */
2321 args->enabled_channels = enabled_channels;
2322
2323 /* Specify whether the EXEC mask represents the valid mask */
2324 args->valid_mask = 0;
2325
2326 /* Specify whether this is the last export */
2327 args->done = 0;
2328
2329 /* Specify the target we are exporting */
2330 args->target = target;
2331
2332 args->compr = false;
2333 args->out[0] = LLVMGetUndef(ctx->ac.f32);
2334 args->out[1] = LLVMGetUndef(ctx->ac.f32);
2335 args->out[2] = LLVMGetUndef(ctx->ac.f32);
2336 args->out[3] = LLVMGetUndef(ctx->ac.f32);
2337
2338 if (!values)
2339 return;
2340
2341 bool is_16bit = ac_get_type_size(LLVMTypeOf(values[0])) == 2;
2342 if (ctx->stage == MESA_SHADER_FRAGMENT) {
2343 unsigned index = target - V_008DFC_SQ_EXP_MRT;
2344 unsigned col_format = (ctx->options->key.fs.col_format >> (4 * index)) & 0xf;
2345 bool is_int8 = (ctx->options->key.fs.is_int8 >> index) & 1;
2346 bool is_int10 = (ctx->options->key.fs.is_int10 >> index) & 1;
2347 unsigned chan;
2348
2349 LLVMValueRef (*packf)(struct ac_llvm_context *ctx, LLVMValueRef args[2]) = NULL;
2350 LLVMValueRef (*packi)(struct ac_llvm_context *ctx, LLVMValueRef args[2],
2351 unsigned bits, bool hi) = NULL;
2352
2353 switch(col_format) {
2354 case V_028714_SPI_SHADER_ZERO:
2355 args->enabled_channels = 0; /* writemask */
2356 args->target = V_008DFC_SQ_EXP_NULL;
2357 break;
2358
2359 case V_028714_SPI_SHADER_32_R:
2360 args->enabled_channels = 1;
2361 args->out[0] = values[0];
2362 break;
2363
2364 case V_028714_SPI_SHADER_32_GR:
2365 args->enabled_channels = 0x3;
2366 args->out[0] = values[0];
2367 args->out[1] = values[1];
2368 break;
2369
2370 case V_028714_SPI_SHADER_32_AR:
2371 if (ctx->ac.chip_class >= GFX10) {
2372 args->enabled_channels = 0x3;
2373 args->out[0] = values[0];
2374 args->out[1] = values[3];
2375 } else {
2376 args->enabled_channels = 0x9;
2377 args->out[0] = values[0];
2378 args->out[3] = values[3];
2379 }
2380 break;
2381
2382 case V_028714_SPI_SHADER_FP16_ABGR:
2383 args->enabled_channels = 0x5;
2384 packf = ac_build_cvt_pkrtz_f16;
2385 if (is_16bit) {
2386 for (unsigned chan = 0; chan < 4; chan++)
2387 values[chan] = LLVMBuildFPExt(ctx->ac.builder,
2388 values[chan],
2389 ctx->ac.f32, "");
2390 }
2391 break;
2392
2393 case V_028714_SPI_SHADER_UNORM16_ABGR:
2394 args->enabled_channels = 0x5;
2395 packf = ac_build_cvt_pknorm_u16;
2396 break;
2397
2398 case V_028714_SPI_SHADER_SNORM16_ABGR:
2399 args->enabled_channels = 0x5;
2400 packf = ac_build_cvt_pknorm_i16;
2401 break;
2402
2403 case V_028714_SPI_SHADER_UINT16_ABGR:
2404 args->enabled_channels = 0x5;
2405 packi = ac_build_cvt_pk_u16;
2406 if (is_16bit) {
2407 for (unsigned chan = 0; chan < 4; chan++)
2408 values[chan] = LLVMBuildZExt(ctx->ac.builder,
2409 ac_to_integer(&ctx->ac, values[chan]),
2410 ctx->ac.i32, "");
2411 }
2412 break;
2413
2414 case V_028714_SPI_SHADER_SINT16_ABGR:
2415 args->enabled_channels = 0x5;
2416 packi = ac_build_cvt_pk_i16;
2417 if (is_16bit) {
2418 for (unsigned chan = 0; chan < 4; chan++)
2419 values[chan] = LLVMBuildSExt(ctx->ac.builder,
2420 ac_to_integer(&ctx->ac, values[chan]),
2421 ctx->ac.i32, "");
2422 }
2423 break;
2424
2425 default:
2426 case V_028714_SPI_SHADER_32_ABGR:
2427 memcpy(&args->out[0], values, sizeof(values[0]) * 4);
2428 break;
2429 }
2430
2431 /* Pack f16 or norm_i16/u16. */
2432 if (packf) {
2433 for (chan = 0; chan < 2; chan++) {
2434 LLVMValueRef pack_args[2] = {
2435 values[2 * chan],
2436 values[2 * chan + 1]
2437 };
2438 LLVMValueRef packed;
2439
2440 packed = packf(&ctx->ac, pack_args);
2441 args->out[chan] = ac_to_float(&ctx->ac, packed);
2442 }
2443 args->compr = 1; /* COMPR flag */
2444 }
2445
2446 /* Pack i16/u16. */
2447 if (packi) {
2448 for (chan = 0; chan < 2; chan++) {
2449 LLVMValueRef pack_args[2] = {
2450 ac_to_integer(&ctx->ac, values[2 * chan]),
2451 ac_to_integer(&ctx->ac, values[2 * chan + 1])
2452 };
2453 LLVMValueRef packed;
2454
2455 packed = packi(&ctx->ac, pack_args,
2456 is_int8 ? 8 : is_int10 ? 10 : 16,
2457 chan == 1);
2458 args->out[chan] = ac_to_float(&ctx->ac, packed);
2459 }
2460 args->compr = 1; /* COMPR flag */
2461 }
2462 return;
2463 }
2464
2465 if (is_16bit) {
2466 for (unsigned chan = 0; chan < 4; chan++) {
2467 values[chan] = LLVMBuildBitCast(ctx->ac.builder, values[chan], ctx->ac.i16, "");
2468 args->out[chan] = LLVMBuildZExt(ctx->ac.builder, values[chan], ctx->ac.i32, "");
2469 }
2470 } else
2471 memcpy(&args->out[0], values, sizeof(values[0]) * 4);
2472
2473 for (unsigned i = 0; i < 4; ++i)
2474 args->out[i] = ac_to_float(&ctx->ac, args->out[i]);
2475 }
2476
2477 static void
2478 radv_export_param(struct radv_shader_context *ctx, unsigned index,
2479 LLVMValueRef *values, unsigned enabled_channels)
2480 {
2481 struct ac_export_args args;
2482
2483 si_llvm_init_export_args(ctx, values, enabled_channels,
2484 V_008DFC_SQ_EXP_PARAM + index, &args);
2485 ac_build_export(&ctx->ac, &args);
2486 }
2487
2488 static LLVMValueRef
2489 radv_load_output(struct radv_shader_context *ctx, unsigned index, unsigned chan)
2490 {
2491 LLVMValueRef output = ctx->abi.outputs[ac_llvm_reg_index_soa(index, chan)];
2492 return LLVMBuildLoad(ctx->ac.builder, output, "");
2493 }
2494
2495 static void
2496 radv_emit_stream_output(struct radv_shader_context *ctx,
2497 LLVMValueRef const *so_buffers,
2498 LLVMValueRef const *so_write_offsets,
2499 const struct radv_stream_output *output,
2500 struct radv_shader_output_values *shader_out)
2501 {
2502 unsigned num_comps = util_bitcount(output->component_mask);
2503 unsigned buf = output->buffer;
2504 unsigned offset = output->offset;
2505 unsigned start;
2506 LLVMValueRef out[4];
2507
2508 assert(num_comps && num_comps <= 4);
2509 if (!num_comps || num_comps > 4)
2510 return;
2511
2512 /* Get the first component. */
2513 start = ffs(output->component_mask) - 1;
2514
2515 /* Load the output as int. */
2516 for (int i = 0; i < num_comps; i++) {
2517 out[i] = ac_to_integer(&ctx->ac, shader_out->values[start + i]);
2518 }
2519
2520 /* Pack the output. */
2521 LLVMValueRef vdata = NULL;
2522
2523 switch (num_comps) {
2524 case 1: /* as i32 */
2525 vdata = out[0];
2526 break;
2527 case 2: /* as v2i32 */
2528 case 3: /* as v4i32 (aligned to 4) */
2529 out[3] = LLVMGetUndef(ctx->ac.i32);
2530 /* fall through */
2531 case 4: /* as v4i32 */
2532 vdata = ac_build_gather_values(&ctx->ac, out,
2533 !ac_has_vec3_support(ctx->ac.chip_class, false) ?
2534 util_next_power_of_two(num_comps) :
2535 num_comps);
2536 break;
2537 }
2538
2539 ac_build_buffer_store_dword(&ctx->ac, so_buffers[buf],
2540 vdata, num_comps, so_write_offsets[buf],
2541 ctx->ac.i32_0, offset,
2542 ac_glc | ac_slc, false);
2543 }
2544
2545 static void
2546 radv_emit_streamout(struct radv_shader_context *ctx, unsigned stream)
2547 {
2548 int i;
2549
2550 /* Get bits [22:16], i.e. (so_param >> 16) & 127; */
2551 assert(ctx->streamout_config);
2552 LLVMValueRef so_vtx_count =
2553 ac_build_bfe(&ctx->ac, ctx->streamout_config,
2554 LLVMConstInt(ctx->ac.i32, 16, false),
2555 LLVMConstInt(ctx->ac.i32, 7, false), false);
2556
2557 LLVMValueRef tid = ac_get_thread_id(&ctx->ac);
2558
2559 /* can_emit = tid < so_vtx_count; */
2560 LLVMValueRef can_emit = LLVMBuildICmp(ctx->ac.builder, LLVMIntULT,
2561 tid, so_vtx_count, "");
2562
2563 /* Emit the streamout code conditionally. This actually avoids
2564 * out-of-bounds buffer access. The hw tells us via the SGPR
2565 * (so_vtx_count) which threads are allowed to emit streamout data.
2566 */
2567 ac_build_ifcc(&ctx->ac, can_emit, 6501);
2568 {
2569 /* The buffer offset is computed as follows:
2570 * ByteOffset = streamout_offset[buffer_id]*4 +
2571 * (streamout_write_index + thread_id)*stride[buffer_id] +
2572 * attrib_offset
2573 */
2574 LLVMValueRef so_write_index = ctx->streamout_write_idx;
2575
2576 /* Compute (streamout_write_index + thread_id). */
2577 so_write_index =
2578 LLVMBuildAdd(ctx->ac.builder, so_write_index, tid, "");
2579
2580 /* Load the descriptor and compute the write offset for each
2581 * enabled buffer.
2582 */
2583 LLVMValueRef so_write_offset[4] = {};
2584 LLVMValueRef so_buffers[4] = {};
2585 LLVMValueRef buf_ptr = ctx->streamout_buffers;
2586
2587 for (i = 0; i < 4; i++) {
2588 uint16_t stride = ctx->shader_info->info.so.strides[i];
2589
2590 if (!stride)
2591 continue;
2592
2593 LLVMValueRef offset =
2594 LLVMConstInt(ctx->ac.i32, i, false);
2595
2596 so_buffers[i] = ac_build_load_to_sgpr(&ctx->ac,
2597 buf_ptr, offset);
2598
2599 LLVMValueRef so_offset = ctx->streamout_offset[i];
2600
2601 so_offset = LLVMBuildMul(ctx->ac.builder, so_offset,
2602 LLVMConstInt(ctx->ac.i32, 4, false), "");
2603
2604 so_write_offset[i] =
2605 ac_build_imad(&ctx->ac, so_write_index,
2606 LLVMConstInt(ctx->ac.i32,
2607 stride * 4, false),
2608 so_offset);
2609 }
2610
2611 /* Write streamout data. */
2612 for (i = 0; i < ctx->shader_info->info.so.num_outputs; i++) {
2613 struct radv_shader_output_values shader_out = {};
2614 struct radv_stream_output *output =
2615 &ctx->shader_info->info.so.outputs[i];
2616
2617 if (stream != output->stream)
2618 continue;
2619
2620 for (int j = 0; j < 4; j++) {
2621 shader_out.values[j] =
2622 radv_load_output(ctx, output->location, j);
2623 }
2624
2625 radv_emit_stream_output(ctx, so_buffers,so_write_offset,
2626 output, &shader_out);
2627 }
2628 }
2629 ac_build_endif(&ctx->ac, 6501);
2630 }
2631
2632 static void
2633 radv_build_param_exports(struct radv_shader_context *ctx,
2634 struct radv_shader_output_values *outputs,
2635 unsigned noutput,
2636 struct radv_vs_output_info *outinfo,
2637 bool export_clip_dists)
2638 {
2639 unsigned param_count = 0;
2640
2641 for (unsigned i = 0; i < noutput; i++) {
2642 unsigned slot_name = outputs[i].slot_name;
2643 unsigned usage_mask = outputs[i].usage_mask;
2644
2645 if (slot_name != VARYING_SLOT_LAYER &&
2646 slot_name != VARYING_SLOT_PRIMITIVE_ID &&
2647 slot_name != VARYING_SLOT_CLIP_DIST0 &&
2648 slot_name != VARYING_SLOT_CLIP_DIST1 &&
2649 slot_name < VARYING_SLOT_VAR0)
2650 continue;
2651
2652 if ((slot_name == VARYING_SLOT_CLIP_DIST0 ||
2653 slot_name == VARYING_SLOT_CLIP_DIST1) && !export_clip_dists)
2654 continue;
2655
2656 radv_export_param(ctx, param_count, outputs[i].values, usage_mask);
2657
2658 assert(i < ARRAY_SIZE(outinfo->vs_output_param_offset));
2659 outinfo->vs_output_param_offset[slot_name] = param_count++;
2660 }
2661
2662 outinfo->param_exports = param_count;
2663 }
2664
2665 /* Generate export instructions for hardware VS shader stage or NGG GS stage
2666 * (position and parameter data only).
2667 */
2668 static void
2669 radv_llvm_export_vs(struct radv_shader_context *ctx,
2670 struct radv_shader_output_values *outputs,
2671 unsigned noutput,
2672 struct radv_vs_output_info *outinfo,
2673 bool export_clip_dists)
2674 {
2675 LLVMValueRef psize_value = NULL, layer_value = NULL, viewport_value = NULL;
2676 struct ac_export_args pos_args[4] = {};
2677 unsigned pos_idx, index;
2678 int i;
2679
2680 /* Build position exports */
2681 for (i = 0; i < noutput; i++) {
2682 switch (outputs[i].slot_name) {
2683 case VARYING_SLOT_POS:
2684 si_llvm_init_export_args(ctx, outputs[i].values, 0xf,
2685 V_008DFC_SQ_EXP_POS, &pos_args[0]);
2686 break;
2687 case VARYING_SLOT_PSIZ:
2688 psize_value = outputs[i].values[0];
2689 break;
2690 case VARYING_SLOT_LAYER:
2691 layer_value = outputs[i].values[0];
2692 break;
2693 case VARYING_SLOT_VIEWPORT:
2694 viewport_value = outputs[i].values[0];
2695 break;
2696 case VARYING_SLOT_CLIP_DIST0:
2697 case VARYING_SLOT_CLIP_DIST1:
2698 index = 2 + outputs[i].slot_index;
2699 si_llvm_init_export_args(ctx, outputs[i].values, 0xf,
2700 V_008DFC_SQ_EXP_POS + index,
2701 &pos_args[index]);
2702 break;
2703 default:
2704 break;
2705 }
2706 }
2707
2708 /* We need to add the position output manually if it's missing. */
2709 if (!pos_args[0].out[0]) {
2710 pos_args[0].enabled_channels = 0xf; /* writemask */
2711 pos_args[0].valid_mask = 0; /* EXEC mask */
2712 pos_args[0].done = 0; /* last export? */
2713 pos_args[0].target = V_008DFC_SQ_EXP_POS;
2714 pos_args[0].compr = 0; /* COMPR flag */
2715 pos_args[0].out[0] = ctx->ac.f32_0; /* X */
2716 pos_args[0].out[1] = ctx->ac.f32_0; /* Y */
2717 pos_args[0].out[2] = ctx->ac.f32_0; /* Z */
2718 pos_args[0].out[3] = ctx->ac.f32_1; /* W */
2719 }
2720
2721 if (outinfo->writes_pointsize ||
2722 outinfo->writes_layer ||
2723 outinfo->writes_viewport_index) {
2724 pos_args[1].enabled_channels = ((outinfo->writes_pointsize == true ? 1 : 0) |
2725 (outinfo->writes_layer == true ? 4 : 0));
2726 pos_args[1].valid_mask = 0;
2727 pos_args[1].done = 0;
2728 pos_args[1].target = V_008DFC_SQ_EXP_POS + 1;
2729 pos_args[1].compr = 0;
2730 pos_args[1].out[0] = ctx->ac.f32_0; /* X */
2731 pos_args[1].out[1] = ctx->ac.f32_0; /* Y */
2732 pos_args[1].out[2] = ctx->ac.f32_0; /* Z */
2733 pos_args[1].out[3] = ctx->ac.f32_0; /* W */
2734
2735 if (outinfo->writes_pointsize == true)
2736 pos_args[1].out[0] = psize_value;
2737 if (outinfo->writes_layer == true)
2738 pos_args[1].out[2] = layer_value;
2739 if (outinfo->writes_viewport_index == true) {
2740 if (ctx->options->chip_class >= GFX9) {
2741 /* GFX9 has the layer in out.z[10:0] and the viewport
2742 * index in out.z[19:16].
2743 */
2744 LLVMValueRef v = viewport_value;
2745 v = ac_to_integer(&ctx->ac, v);
2746 v = LLVMBuildShl(ctx->ac.builder, v,
2747 LLVMConstInt(ctx->ac.i32, 16, false),
2748 "");
2749 v = LLVMBuildOr(ctx->ac.builder, v,
2750 ac_to_integer(&ctx->ac, pos_args[1].out[2]), "");
2751
2752 pos_args[1].out[2] = ac_to_float(&ctx->ac, v);
2753 pos_args[1].enabled_channels |= 1 << 2;
2754 } else {
2755 pos_args[1].out[3] = viewport_value;
2756 pos_args[1].enabled_channels |= 1 << 3;
2757 }
2758 }
2759 }
2760
2761 for (i = 0; i < 4; i++) {
2762 if (pos_args[i].out[0])
2763 outinfo->pos_exports++;
2764 }
2765
2766 /* Navi10-14 skip POS0 exports if EXEC=0 and DONE=0, causing a hang.
2767 * Setting valid_mask=1 prevents it and has no other effect.
2768 */
2769 if (ctx->ac.family == CHIP_NAVI10 ||
2770 ctx->ac.family == CHIP_NAVI12 ||
2771 ctx->ac.family == CHIP_NAVI14)
2772 pos_args[0].valid_mask = 1;
2773
2774 pos_idx = 0;
2775 for (i = 0; i < 4; i++) {
2776 if (!pos_args[i].out[0])
2777 continue;
2778
2779 /* Specify the target we are exporting */
2780 pos_args[i].target = V_008DFC_SQ_EXP_POS + pos_idx++;
2781
2782 if (pos_idx == outinfo->pos_exports)
2783 /* Specify that this is the last export */
2784 pos_args[i].done = 1;
2785
2786 ac_build_export(&ctx->ac, &pos_args[i]);
2787 }
2788
2789 /* Build parameter exports */
2790 radv_build_param_exports(ctx, outputs, noutput, outinfo, export_clip_dists);
2791 }
2792
2793 static void
2794 handle_vs_outputs_post(struct radv_shader_context *ctx,
2795 bool export_prim_id,
2796 bool export_clip_dists,
2797 struct radv_vs_output_info *outinfo)
2798 {
2799 struct radv_shader_output_values *outputs;
2800 unsigned noutput = 0;
2801
2802 if (ctx->options->key.has_multiview_view_index) {
2803 LLVMValueRef* tmp_out = &ctx->abi.outputs[ac_llvm_reg_index_soa(VARYING_SLOT_LAYER, 0)];
2804 if(!*tmp_out) {
2805 for(unsigned i = 0; i < 4; ++i)
2806 ctx->abi.outputs[ac_llvm_reg_index_soa(VARYING_SLOT_LAYER, i)] =
2807 ac_build_alloca_undef(&ctx->ac, ctx->ac.f32, "");
2808 }
2809
2810 LLVMBuildStore(ctx->ac.builder, ac_to_float(&ctx->ac, ctx->abi.view_index), *tmp_out);
2811 ctx->output_mask |= 1ull << VARYING_SLOT_LAYER;
2812 }
2813
2814 memset(outinfo->vs_output_param_offset, AC_EXP_PARAM_UNDEFINED,
2815 sizeof(outinfo->vs_output_param_offset));
2816 outinfo->pos_exports = 0;
2817
2818 if (ctx->output_mask & (1ull << VARYING_SLOT_PSIZ)) {
2819 outinfo->writes_pointsize = true;
2820 }
2821
2822 if (ctx->output_mask & (1ull << VARYING_SLOT_LAYER)) {
2823 outinfo->writes_layer = true;
2824 }
2825
2826 if (ctx->output_mask & (1ull << VARYING_SLOT_VIEWPORT)) {
2827 outinfo->writes_viewport_index = true;
2828 }
2829
2830 if (ctx->shader_info->info.so.num_outputs &&
2831 !ctx->is_gs_copy_shader) {
2832 /* The GS copy shader emission already emits streamout. */
2833 radv_emit_streamout(ctx, 0);
2834 }
2835
2836 /* Allocate a temporary array for the output values. */
2837 unsigned num_outputs = util_bitcount64(ctx->output_mask) + export_prim_id;
2838 outputs = malloc(num_outputs * sizeof(outputs[0]));
2839
2840 for (unsigned i = 0; i < AC_LLVM_MAX_OUTPUTS; ++i) {
2841 if (!(ctx->output_mask & (1ull << i)))
2842 continue;
2843
2844 outputs[noutput].slot_name = i;
2845 outputs[noutput].slot_index = i == VARYING_SLOT_CLIP_DIST1;
2846
2847 if (ctx->stage == MESA_SHADER_VERTEX &&
2848 !ctx->is_gs_copy_shader) {
2849 outputs[noutput].usage_mask =
2850 ctx->shader_info->info.vs.output_usage_mask[i];
2851 } else if (ctx->stage == MESA_SHADER_TESS_EVAL) {
2852 outputs[noutput].usage_mask =
2853 ctx->shader_info->info.tes.output_usage_mask[i];
2854 } else {
2855 assert(ctx->is_gs_copy_shader);
2856 outputs[noutput].usage_mask =
2857 ctx->shader_info->info.gs.output_usage_mask[i];
2858 }
2859
2860 for (unsigned j = 0; j < 4; j++) {
2861 outputs[noutput].values[j] =
2862 ac_to_float(&ctx->ac, radv_load_output(ctx, i, j));
2863 }
2864
2865 noutput++;
2866 }
2867
2868 /* Export PrimitiveID. */
2869 if (export_prim_id) {
2870 outinfo->export_prim_id = true;
2871
2872 outputs[noutput].slot_name = VARYING_SLOT_PRIMITIVE_ID;
2873 outputs[noutput].slot_index = 0;
2874 outputs[noutput].usage_mask = 0x1;
2875 outputs[noutput].values[0] = ctx->vs_prim_id;
2876 for (unsigned j = 1; j < 4; j++)
2877 outputs[noutput].values[j] = ctx->ac.f32_0;
2878 noutput++;
2879 }
2880
2881 radv_llvm_export_vs(ctx, outputs, noutput, outinfo, export_clip_dists);
2882
2883 free(outputs);
2884 }
2885
2886 static void
2887 handle_es_outputs_post(struct radv_shader_context *ctx,
2888 struct radv_es_output_info *outinfo)
2889 {
2890 int j;
2891 uint64_t max_output_written = 0;
2892 LLVMValueRef lds_base = NULL;
2893
2894 for (unsigned i = 0; i < AC_LLVM_MAX_OUTPUTS; ++i) {
2895 int param_index;
2896
2897 if (!(ctx->output_mask & (1ull << i)))
2898 continue;
2899
2900 param_index = shader_io_get_unique_index(i);
2901
2902 max_output_written = MAX2(param_index, max_output_written);
2903 }
2904
2905 outinfo->esgs_itemsize = (max_output_written + 1) * 16;
2906
2907 if (ctx->ac.chip_class >= GFX9) {
2908 unsigned itemsize_dw = outinfo->esgs_itemsize / 4;
2909 LLVMValueRef vertex_idx = ac_get_thread_id(&ctx->ac);
2910 LLVMValueRef wave_idx = ac_unpack_param(&ctx->ac, ctx->merged_wave_info, 24, 4);
2911 vertex_idx = LLVMBuildOr(ctx->ac.builder, vertex_idx,
2912 LLVMBuildMul(ctx->ac.builder, wave_idx,
2913 LLVMConstInt(ctx->ac.i32,
2914 ctx->ac.wave_size, false), ""), "");
2915 lds_base = LLVMBuildMul(ctx->ac.builder, vertex_idx,
2916 LLVMConstInt(ctx->ac.i32, itemsize_dw, 0), "");
2917 }
2918
2919 for (unsigned i = 0; i < AC_LLVM_MAX_OUTPUTS; ++i) {
2920 LLVMValueRef dw_addr = NULL;
2921 LLVMValueRef *out_ptr = &ctx->abi.outputs[i * 4];
2922 unsigned output_usage_mask;
2923 int param_index;
2924
2925 if (!(ctx->output_mask & (1ull << i)))
2926 continue;
2927
2928 if (ctx->stage == MESA_SHADER_VERTEX) {
2929 output_usage_mask =
2930 ctx->shader_info->info.vs.output_usage_mask[i];
2931 } else {
2932 assert(ctx->stage == MESA_SHADER_TESS_EVAL);
2933 output_usage_mask =
2934 ctx->shader_info->info.tes.output_usage_mask[i];
2935 }
2936
2937 param_index = shader_io_get_unique_index(i);
2938
2939 if (lds_base) {
2940 dw_addr = LLVMBuildAdd(ctx->ac.builder, lds_base,
2941 LLVMConstInt(ctx->ac.i32, param_index * 4, false),
2942 "");
2943 }
2944
2945 for (j = 0; j < 4; j++) {
2946 if (!(output_usage_mask & (1 << j)))
2947 continue;
2948
2949 LLVMValueRef out_val = LLVMBuildLoad(ctx->ac.builder, out_ptr[j], "");
2950 out_val = ac_to_integer(&ctx->ac, out_val);
2951 out_val = LLVMBuildZExtOrBitCast(ctx->ac.builder, out_val, ctx->ac.i32, "");
2952
2953 if (ctx->ac.chip_class >= GFX9) {
2954 LLVMValueRef dw_addr_offset =
2955 LLVMBuildAdd(ctx->ac.builder, dw_addr,
2956 LLVMConstInt(ctx->ac.i32,
2957 j, false), "");
2958
2959 ac_lds_store(&ctx->ac, dw_addr_offset, out_val);
2960 } else {
2961 ac_build_buffer_store_dword(&ctx->ac,
2962 ctx->esgs_ring,
2963 out_val, 1,
2964 NULL, ctx->es2gs_offset,
2965 (4 * param_index + j) * 4,
2966 ac_glc | ac_slc, true);
2967 }
2968 }
2969 }
2970 }
2971
2972 static void
2973 handle_ls_outputs_post(struct radv_shader_context *ctx)
2974 {
2975 LLVMValueRef vertex_id = ctx->rel_auto_id;
2976 uint32_t num_tcs_inputs = util_last_bit64(ctx->shader_info->info.vs.ls_outputs_written);
2977 LLVMValueRef vertex_dw_stride = LLVMConstInt(ctx->ac.i32, num_tcs_inputs * 4, false);
2978 LLVMValueRef base_dw_addr = LLVMBuildMul(ctx->ac.builder, vertex_id,
2979 vertex_dw_stride, "");
2980
2981 for (unsigned i = 0; i < AC_LLVM_MAX_OUTPUTS; ++i) {
2982 LLVMValueRef *out_ptr = &ctx->abi.outputs[i * 4];
2983
2984 if (!(ctx->output_mask & (1ull << i)))
2985 continue;
2986
2987 int param = shader_io_get_unique_index(i);
2988 LLVMValueRef dw_addr = LLVMBuildAdd(ctx->ac.builder, base_dw_addr,
2989 LLVMConstInt(ctx->ac.i32, param * 4, false),
2990 "");
2991 for (unsigned j = 0; j < 4; j++) {
2992 LLVMValueRef value = LLVMBuildLoad(ctx->ac.builder, out_ptr[j], "");
2993 value = ac_to_integer(&ctx->ac, value);
2994 value = LLVMBuildZExtOrBitCast(ctx->ac.builder, value, ctx->ac.i32, "");
2995 ac_lds_store(&ctx->ac, dw_addr, value);
2996 dw_addr = LLVMBuildAdd(ctx->ac.builder, dw_addr, ctx->ac.i32_1, "");
2997 }
2998 }
2999 }
3000
3001 static LLVMValueRef get_wave_id_in_tg(struct radv_shader_context *ctx)
3002 {
3003 return ac_unpack_param(&ctx->ac, ctx->merged_wave_info, 24, 4);
3004 }
3005
3006 static LLVMValueRef get_tgsize(struct radv_shader_context *ctx)
3007 {
3008 return ac_unpack_param(&ctx->ac, ctx->merged_wave_info, 28, 4);
3009 }
3010
3011 static LLVMValueRef get_thread_id_in_tg(struct radv_shader_context *ctx)
3012 {
3013 LLVMBuilderRef builder = ctx->ac.builder;
3014 LLVMValueRef tmp;
3015 tmp = LLVMBuildMul(builder, get_wave_id_in_tg(ctx),
3016 LLVMConstInt(ctx->ac.i32, ctx->ac.wave_size, false), "");
3017 return LLVMBuildAdd(builder, tmp, ac_get_thread_id(&ctx->ac), "");
3018 }
3019
3020 static LLVMValueRef ngg_get_vtx_cnt(struct radv_shader_context *ctx)
3021 {
3022 return ac_build_bfe(&ctx->ac, ctx->gs_tg_info,
3023 LLVMConstInt(ctx->ac.i32, 12, false),
3024 LLVMConstInt(ctx->ac.i32, 9, false),
3025 false);
3026 }
3027
3028 static LLVMValueRef ngg_get_prim_cnt(struct radv_shader_context *ctx)
3029 {
3030 return ac_build_bfe(&ctx->ac, ctx->gs_tg_info,
3031 LLVMConstInt(ctx->ac.i32, 22, false),
3032 LLVMConstInt(ctx->ac.i32, 9, false),
3033 false);
3034 }
3035
3036 static LLVMValueRef
3037 ngg_gs_get_vertex_storage(struct radv_shader_context *ctx)
3038 {
3039 unsigned num_outputs = util_bitcount64(ctx->output_mask);
3040
3041 LLVMTypeRef elements[2] = {
3042 LLVMArrayType(ctx->ac.i32, 4 * num_outputs),
3043 LLVMArrayType(ctx->ac.i8, 4),
3044 };
3045 LLVMTypeRef type = LLVMStructTypeInContext(ctx->ac.context, elements, 2, false);
3046 type = LLVMPointerType(LLVMArrayType(type, 0), AC_ADDR_SPACE_LDS);
3047 return LLVMBuildBitCast(ctx->ac.builder, ctx->gs_ngg_emit, type, "");
3048 }
3049
3050 /**
3051 * Return a pointer to the LDS storage reserved for the N'th vertex, where N
3052 * is in emit order; that is:
3053 * - during the epilogue, N is the threadidx (relative to the entire threadgroup)
3054 * - during vertex emit, i.e. while the API GS shader invocation is running,
3055 * N = threadidx * gs_max_out_vertices + emitidx
3056 *
3057 * Goals of the LDS memory layout:
3058 * 1. Eliminate bank conflicts on write for geometry shaders that have all emits
3059 * in uniform control flow
3060 * 2. Eliminate bank conflicts on read for export if, additionally, there is no
3061 * culling
3062 * 3. Agnostic to the number of waves (since we don't know it before compiling)
3063 * 4. Allow coalescing of LDS instructions (ds_write_b128 etc.)
3064 * 5. Avoid wasting memory.
3065 *
3066 * We use an AoS layout due to point 4 (this also helps point 3). In an AoS
3067 * layout, elimination of bank conflicts requires that each vertex occupy an
3068 * odd number of dwords. We use the additional dword to store the output stream
3069 * index as well as a flag to indicate whether this vertex ends a primitive
3070 * for rasterization.
3071 *
3072 * Swizzling is required to satisfy points 1 and 2 simultaneously.
3073 *
3074 * Vertices are stored in export order (gsthread * gs_max_out_vertices + emitidx).
3075 * Indices are swizzled in groups of 32, which ensures point 1 without
3076 * disturbing point 2.
3077 *
3078 * \return an LDS pointer to type {[N x i32], [4 x i8]}
3079 */
3080 static LLVMValueRef
3081 ngg_gs_vertex_ptr(struct radv_shader_context *ctx, LLVMValueRef vertexidx)
3082 {
3083 LLVMBuilderRef builder = ctx->ac.builder;
3084 LLVMValueRef storage = ngg_gs_get_vertex_storage(ctx);
3085
3086 /* gs_max_out_vertices = 2^(write_stride_2exp) * some odd number */
3087 unsigned write_stride_2exp = ffs(ctx->gs_max_out_vertices) - 1;
3088 if (write_stride_2exp) {
3089 LLVMValueRef row =
3090 LLVMBuildLShr(builder, vertexidx,
3091 LLVMConstInt(ctx->ac.i32, 5, false), "");
3092 LLVMValueRef swizzle =
3093 LLVMBuildAnd(builder, row,
3094 LLVMConstInt(ctx->ac.i32, (1u << write_stride_2exp) - 1,
3095 false), "");
3096 vertexidx = LLVMBuildXor(builder, vertexidx, swizzle, "");
3097 }
3098
3099 return ac_build_gep0(&ctx->ac, storage, vertexidx);
3100 }
3101
3102 static LLVMValueRef
3103 ngg_gs_emit_vertex_ptr(struct radv_shader_context *ctx, LLVMValueRef gsthread,
3104 LLVMValueRef emitidx)
3105 {
3106 LLVMBuilderRef builder = ctx->ac.builder;
3107 LLVMValueRef tmp;
3108
3109 tmp = LLVMConstInt(ctx->ac.i32, ctx->gs_max_out_vertices, false);
3110 tmp = LLVMBuildMul(builder, tmp, gsthread, "");
3111 const LLVMValueRef vertexidx = LLVMBuildAdd(builder, tmp, emitidx, "");
3112 return ngg_gs_vertex_ptr(ctx, vertexidx);
3113 }
3114
3115 /* Send GS Alloc Req message from the first wave of the group to SPI.
3116 * Message payload is:
3117 * - bits 0..10: vertices in group
3118 * - bits 12..22: primitives in group
3119 */
3120 static void build_sendmsg_gs_alloc_req(struct radv_shader_context *ctx,
3121 LLVMValueRef vtx_cnt,
3122 LLVMValueRef prim_cnt)
3123 {
3124 LLVMBuilderRef builder = ctx->ac.builder;
3125 LLVMValueRef tmp;
3126
3127 tmp = LLVMBuildICmp(builder, LLVMIntEQ, get_wave_id_in_tg(ctx), ctx->ac.i32_0, "");
3128 ac_build_ifcc(&ctx->ac, tmp, 5020);
3129
3130 tmp = LLVMBuildShl(builder, prim_cnt, LLVMConstInt(ctx->ac.i32, 12, false),"");
3131 tmp = LLVMBuildOr(builder, tmp, vtx_cnt, "");
3132 ac_build_sendmsg(&ctx->ac, AC_SENDMSG_GS_ALLOC_REQ, tmp);
3133
3134 ac_build_endif(&ctx->ac, 5020);
3135 }
3136
3137 struct ngg_prim {
3138 unsigned num_vertices;
3139 LLVMValueRef isnull;
3140 LLVMValueRef index[3];
3141 LLVMValueRef edgeflag[3];
3142 };
3143
3144 static void build_export_prim(struct radv_shader_context *ctx,
3145 const struct ngg_prim *prim)
3146 {
3147 LLVMBuilderRef builder = ctx->ac.builder;
3148 struct ac_export_args args;
3149 LLVMValueRef tmp;
3150
3151 tmp = LLVMBuildZExt(builder, prim->isnull, ctx->ac.i32, "");
3152 args.out[0] = LLVMBuildShl(builder, tmp, LLVMConstInt(ctx->ac.i32, 31, false), "");
3153
3154 for (unsigned i = 0; i < prim->num_vertices; ++i) {
3155 tmp = LLVMBuildShl(builder, prim->index[i],
3156 LLVMConstInt(ctx->ac.i32, 10 * i, false), "");
3157 args.out[0] = LLVMBuildOr(builder, args.out[0], tmp, "");
3158 tmp = LLVMBuildZExt(builder, prim->edgeflag[i], ctx->ac.i32, "");
3159 tmp = LLVMBuildShl(builder, tmp,
3160 LLVMConstInt(ctx->ac.i32, 10 * i + 9, false), "");
3161 args.out[0] = LLVMBuildOr(builder, args.out[0], tmp, "");
3162 }
3163
3164 args.out[0] = LLVMBuildBitCast(builder, args.out[0], ctx->ac.f32, "");
3165 args.out[1] = LLVMGetUndef(ctx->ac.f32);
3166 args.out[2] = LLVMGetUndef(ctx->ac.f32);
3167 args.out[3] = LLVMGetUndef(ctx->ac.f32);
3168
3169 args.target = V_008DFC_SQ_EXP_PRIM;
3170 args.enabled_channels = 1;
3171 args.done = true;
3172 args.valid_mask = false;
3173 args.compr = false;
3174
3175 ac_build_export(&ctx->ac, &args);
3176 }
3177
3178 static void
3179 handle_ngg_outputs_post(struct radv_shader_context *ctx)
3180 {
3181 LLVMBuilderRef builder = ctx->ac.builder;
3182 unsigned num_vertices = 3;
3183 LLVMValueRef tmp;
3184
3185 assert((ctx->stage == MESA_SHADER_VERTEX ||
3186 ctx->stage == MESA_SHADER_TESS_EVAL) && !ctx->is_gs_copy_shader);
3187
3188 LLVMValueRef prims_in_wave = ac_unpack_param(&ctx->ac, ctx->merged_wave_info, 8, 8);
3189 LLVMValueRef vtx_in_wave = ac_unpack_param(&ctx->ac, ctx->merged_wave_info, 0, 8);
3190 LLVMValueRef is_gs_thread = LLVMBuildICmp(builder, LLVMIntULT,
3191 ac_get_thread_id(&ctx->ac), prims_in_wave, "");
3192 LLVMValueRef is_es_thread = LLVMBuildICmp(builder, LLVMIntULT,
3193 ac_get_thread_id(&ctx->ac), vtx_in_wave, "");
3194 LLVMValueRef vtxindex[] = {
3195 ac_unpack_param(&ctx->ac, ctx->gs_vtx_offset[0], 0, 16),
3196 ac_unpack_param(&ctx->ac, ctx->gs_vtx_offset[0], 16, 16),
3197 ac_unpack_param(&ctx->ac, ctx->gs_vtx_offset[2], 0, 16),
3198 };
3199
3200 /* TODO: streamout */
3201
3202 /* Copy Primitive IDs from GS threads to the LDS address corresponding
3203 * to the ES thread of the provoking vertex.
3204 */
3205 if (ctx->stage == MESA_SHADER_VERTEX &&
3206 ctx->options->key.vs_common_out.export_prim_id) {
3207 /* TODO: streamout */
3208
3209 ac_build_ifcc(&ctx->ac, is_gs_thread, 5400);
3210 /* Extract the PROVOKING_VTX_INDEX field. */
3211 LLVMValueRef provoking_vtx_in_prim =
3212 LLVMConstInt(ctx->ac.i32, 0, false);
3213
3214 /* provoking_vtx_index = vtxindex[provoking_vtx_in_prim]; */
3215 LLVMValueRef indices = ac_build_gather_values(&ctx->ac, vtxindex, 3);
3216 LLVMValueRef provoking_vtx_index =
3217 LLVMBuildExtractElement(builder, indices, provoking_vtx_in_prim, "");
3218
3219 LLVMBuildStore(builder, ctx->abi.gs_prim_id,
3220 ac_build_gep0(&ctx->ac, ctx->esgs_ring, provoking_vtx_index));
3221 ac_build_endif(&ctx->ac, 5400);
3222 }
3223
3224 /* TODO: primitive culling */
3225
3226 build_sendmsg_gs_alloc_req(ctx, ngg_get_vtx_cnt(ctx), ngg_get_prim_cnt(ctx));
3227
3228 /* TODO: streamout queries */
3229 /* Export primitive data to the index buffer. Format is:
3230 * - bits 0..8: index 0
3231 * - bit 9: edge flag 0
3232 * - bits 10..18: index 1
3233 * - bit 19: edge flag 1
3234 * - bits 20..28: index 2
3235 * - bit 29: edge flag 2
3236 * - bit 31: null primitive (skip)
3237 *
3238 * For the first version, we will always build up all three indices
3239 * independent of the primitive type. The additional garbage data
3240 * shouldn't hurt.
3241 *
3242 * TODO: culling depends on the primitive type, so can have some
3243 * interaction here.
3244 */
3245 ac_build_ifcc(&ctx->ac, is_gs_thread, 6001);
3246 {
3247 struct ngg_prim prim = {};
3248
3249 prim.num_vertices = num_vertices;
3250 prim.isnull = ctx->ac.i1false;
3251 memcpy(prim.index, vtxindex, sizeof(vtxindex[0]) * 3);
3252
3253 for (unsigned i = 0; i < num_vertices; ++i) {
3254 tmp = LLVMBuildLShr(builder, ctx->abi.gs_invocation_id,
3255 LLVMConstInt(ctx->ac.i32, 8 + i, false), "");
3256 prim.edgeflag[i] = LLVMBuildTrunc(builder, tmp, ctx->ac.i1, "");
3257 }
3258
3259 build_export_prim(ctx, &prim);
3260 }
3261 ac_build_endif(&ctx->ac, 6001);
3262
3263 /* Export per-vertex data (positions and parameters). */
3264 ac_build_ifcc(&ctx->ac, is_es_thread, 6002);
3265 {
3266 struct radv_vs_output_info *outinfo =
3267 ctx->stage == MESA_SHADER_TESS_EVAL ? &ctx->shader_info->tes.outinfo : &ctx->shader_info->vs.outinfo;
3268
3269 /* Exporting the primitive ID is handled below. */
3270 /* TODO: use the new VS export path */
3271 handle_vs_outputs_post(ctx, false,
3272 ctx->options->key.vs_common_out.export_clip_dists,
3273 outinfo);
3274
3275 if (ctx->options->key.vs_common_out.export_prim_id) {
3276 unsigned param_count = outinfo->param_exports;
3277 LLVMValueRef values[4];
3278
3279 if (ctx->stage == MESA_SHADER_VERTEX) {
3280 /* Wait for GS stores to finish. */
3281 ac_build_s_barrier(&ctx->ac);
3282
3283 tmp = ac_build_gep0(&ctx->ac, ctx->esgs_ring,
3284 get_thread_id_in_tg(ctx));
3285 values[0] = LLVMBuildLoad(builder, tmp, "");
3286 } else {
3287 assert(ctx->stage == MESA_SHADER_TESS_EVAL);
3288 values[0] = ctx->abi.tes_patch_id;
3289 }
3290
3291 values[0] = ac_to_float(&ctx->ac, values[0]);
3292 for (unsigned j = 1; j < 4; j++)
3293 values[j] = ctx->ac.f32_0;
3294
3295 radv_export_param(ctx, param_count, values, 0x1);
3296
3297 outinfo->vs_output_param_offset[VARYING_SLOT_PRIMITIVE_ID] = param_count++;
3298 outinfo->export_prim_id = true;
3299 outinfo->param_exports = param_count;
3300 }
3301 }
3302 ac_build_endif(&ctx->ac, 6002);
3303 }
3304
3305 static void gfx10_ngg_gs_emit_prologue(struct radv_shader_context *ctx)
3306 {
3307 /* Zero out the part of LDS scratch that is used to accumulate the
3308 * per-stream generated primitive count.
3309 */
3310 LLVMBuilderRef builder = ctx->ac.builder;
3311 LLVMValueRef scratchptr = ctx->gs_ngg_scratch;
3312 LLVMValueRef tid = get_thread_id_in_tg(ctx);
3313 LLVMBasicBlockRef merge_block;
3314 LLVMValueRef cond;
3315
3316 LLVMValueRef fn = LLVMGetBasicBlockParent(LLVMGetInsertBlock(ctx->ac.builder));
3317 LLVMBasicBlockRef then_block = LLVMAppendBasicBlockInContext(ctx->ac.context, fn, "");
3318 merge_block = LLVMAppendBasicBlockInContext(ctx->ac.context, fn, "");
3319
3320 cond = LLVMBuildICmp(builder, LLVMIntULT, tid, LLVMConstInt(ctx->ac.i32, 4, false), "");
3321 LLVMBuildCondBr(ctx->ac.builder, cond, then_block, merge_block);
3322 LLVMPositionBuilderAtEnd(ctx->ac.builder, then_block);
3323
3324 LLVMValueRef ptr = ac_build_gep0(&ctx->ac, scratchptr, tid);
3325 LLVMBuildStore(builder, ctx->ac.i32_0, ptr);
3326
3327 LLVMBuildBr(ctx->ac.builder, merge_block);
3328 LLVMPositionBuilderAtEnd(ctx->ac.builder, merge_block);
3329
3330 ac_build_s_barrier(&ctx->ac);
3331 }
3332
3333 static void gfx10_ngg_gs_emit_epilogue_1(struct radv_shader_context *ctx)
3334 {
3335 LLVMBuilderRef builder = ctx->ac.builder;
3336 LLVMValueRef i8_0 = LLVMConstInt(ctx->ac.i8, 0, false);
3337 LLVMValueRef tmp;
3338
3339 /* Zero out remaining (non-emitted) primitive flags.
3340 *
3341 * Note: Alternatively, we could pass the relevant gs_next_vertex to
3342 * the emit threads via LDS. This is likely worse in the expected
3343 * typical case where each GS thread emits the full set of
3344 * vertices.
3345 */
3346 for (unsigned stream = 0; stream < 4; ++stream) {
3347 unsigned num_components;
3348
3349 num_components =
3350 ctx->shader_info->info.gs.num_stream_output_components[stream];
3351 if (!num_components)
3352 continue;
3353
3354 const LLVMValueRef gsthread = get_thread_id_in_tg(ctx);
3355
3356 ac_build_bgnloop(&ctx->ac, 5100);
3357
3358 const LLVMValueRef vertexidx =
3359 LLVMBuildLoad(builder, ctx->gs_next_vertex[stream], "");
3360 tmp = LLVMBuildICmp(builder, LLVMIntUGE, vertexidx,
3361 LLVMConstInt(ctx->ac.i32, ctx->gs_max_out_vertices, false), "");
3362 ac_build_ifcc(&ctx->ac, tmp, 5101);
3363 ac_build_break(&ctx->ac);
3364 ac_build_endif(&ctx->ac, 5101);
3365
3366 tmp = LLVMBuildAdd(builder, vertexidx, ctx->ac.i32_1, "");
3367 LLVMBuildStore(builder, tmp, ctx->gs_next_vertex[stream]);
3368
3369 tmp = ngg_gs_emit_vertex_ptr(ctx, gsthread, vertexidx);
3370 LLVMValueRef gep_idx[3] = {
3371 ctx->ac.i32_0, /* implied C-style array */
3372 ctx->ac.i32_1, /* second entry of struct */
3373 LLVMConstInt(ctx->ac.i32, stream, false),
3374 };
3375 tmp = LLVMBuildGEP(builder, tmp, gep_idx, 3, "");
3376 LLVMBuildStore(builder, i8_0, tmp);
3377
3378 ac_build_endloop(&ctx->ac, 5100);
3379 }
3380 }
3381
3382 static void gfx10_ngg_gs_emit_epilogue_2(struct radv_shader_context *ctx)
3383 {
3384 const unsigned verts_per_prim = si_conv_gl_prim_to_vertices(ctx->gs_output_prim);
3385 LLVMBuilderRef builder = ctx->ac.builder;
3386 LLVMValueRef tmp, tmp2;
3387
3388 ac_build_s_barrier(&ctx->ac);
3389
3390 const LLVMValueRef tid = get_thread_id_in_tg(ctx);
3391 LLVMValueRef num_emit_threads = ngg_get_prim_cnt(ctx);
3392
3393 /* TODO: streamout */
3394
3395 /* TODO: culling */
3396
3397 /* Determine vertex liveness. */
3398 LLVMValueRef vertliveptr = ac_build_alloca(&ctx->ac, ctx->ac.i1, "vertexlive");
3399
3400 tmp = LLVMBuildICmp(builder, LLVMIntULT, tid, num_emit_threads, "");
3401 ac_build_ifcc(&ctx->ac, tmp, 5120);
3402 {
3403 for (unsigned i = 0; i < verts_per_prim; ++i) {
3404 const LLVMValueRef primidx =
3405 LLVMBuildAdd(builder, tid,
3406 LLVMConstInt(ctx->ac.i32, i, false), "");
3407
3408 if (i > 0) {
3409 tmp = LLVMBuildICmp(builder, LLVMIntULT, primidx, num_emit_threads, "");
3410 ac_build_ifcc(&ctx->ac, tmp, 5121 + i);
3411 }
3412
3413 /* Load primitive liveness */
3414 tmp = ngg_gs_vertex_ptr(ctx, primidx);
3415 LLVMValueRef gep_idx[3] = {
3416 ctx->ac.i32_0, /* implicit C-style array */
3417 ctx->ac.i32_1, /* second value of struct */
3418 ctx->ac.i32_0, /* stream 0 */
3419 };
3420 tmp = LLVMBuildGEP(builder, tmp, gep_idx, 3, "");
3421 tmp = LLVMBuildLoad(builder, tmp, "");
3422 const LLVMValueRef primlive =
3423 LLVMBuildTrunc(builder, tmp, ctx->ac.i1, "");
3424
3425 tmp = LLVMBuildLoad(builder, vertliveptr, "");
3426 tmp = LLVMBuildOr(builder, tmp, primlive, ""),
3427 LLVMBuildStore(builder, tmp, vertliveptr);
3428
3429 if (i > 0)
3430 ac_build_endif(&ctx->ac, 5121 + i);
3431 }
3432 }
3433 ac_build_endif(&ctx->ac, 5120);
3434
3435 /* Inclusive scan addition across the current wave. */
3436 LLVMValueRef vertlive = LLVMBuildLoad(builder, vertliveptr, "");
3437 struct ac_wg_scan vertlive_scan = {};
3438 vertlive_scan.op = nir_op_iadd;
3439 vertlive_scan.enable_reduce = true;
3440 vertlive_scan.enable_exclusive = true;
3441 vertlive_scan.src = vertlive;
3442 vertlive_scan.scratch = ac_build_gep0(&ctx->ac, ctx->gs_ngg_scratch, ctx->ac.i32_0);
3443 vertlive_scan.waveidx = get_wave_id_in_tg(ctx);
3444 vertlive_scan.numwaves = get_tgsize(ctx);
3445 vertlive_scan.maxwaves = 8;
3446
3447 ac_build_wg_scan(&ctx->ac, &vertlive_scan);
3448
3449 /* Skip all exports (including index exports) when possible. At least on
3450 * early gfx10 revisions this is also to avoid hangs.
3451 */
3452 LLVMValueRef have_exports =
3453 LLVMBuildICmp(builder, LLVMIntNE, vertlive_scan.result_reduce, ctx->ac.i32_0, "");
3454 num_emit_threads =
3455 LLVMBuildSelect(builder, have_exports, num_emit_threads, ctx->ac.i32_0, "");
3456
3457 /* Allocate export space. Send this message as early as possible, to
3458 * hide the latency of the SQ <-> SPI roundtrip.
3459 *
3460 * Note: We could consider compacting primitives for export as well.
3461 * PA processes 1 non-null prim / clock, but it fetches 4 DW of
3462 * prim data per clock and skips null primitives at no additional
3463 * cost. So compacting primitives can only be beneficial when
3464 * there are 4 or more contiguous null primitives in the export
3465 * (in the common case of single-dword prim exports).
3466 */
3467 build_sendmsg_gs_alloc_req(ctx, vertlive_scan.result_reduce, num_emit_threads);
3468
3469 /* Setup the reverse vertex compaction permutation. We re-use stream 1
3470 * of the primitive liveness flags, relying on the fact that each
3471 * threadgroup can have at most 256 threads. */
3472 ac_build_ifcc(&ctx->ac, vertlive, 5130);
3473 {
3474 tmp = ngg_gs_vertex_ptr(ctx, vertlive_scan.result_exclusive);
3475 LLVMValueRef gep_idx[3] = {
3476 ctx->ac.i32_0, /* implicit C-style array */
3477 ctx->ac.i32_1, /* second value of struct */
3478 ctx->ac.i32_1, /* stream 1 */
3479 };
3480 tmp = LLVMBuildGEP(builder, tmp, gep_idx, 3, "");
3481 tmp2 = LLVMBuildTrunc(builder, tid, ctx->ac.i8, "");
3482 LLVMBuildStore(builder, tmp2, tmp);
3483 }
3484 ac_build_endif(&ctx->ac, 5130);
3485
3486 ac_build_s_barrier(&ctx->ac);
3487
3488 /* Export primitive data */
3489 tmp = LLVMBuildICmp(builder, LLVMIntULT, tid, num_emit_threads, "");
3490 ac_build_ifcc(&ctx->ac, tmp, 5140);
3491 {
3492 struct ngg_prim prim = {};
3493 prim.num_vertices = verts_per_prim;
3494
3495 tmp = ngg_gs_vertex_ptr(ctx, tid);
3496 LLVMValueRef gep_idx[3] = {
3497 ctx->ac.i32_0, /* implicit C-style array */
3498 ctx->ac.i32_1, /* second value of struct */
3499 ctx->ac.i32_0, /* primflag */
3500 };
3501 tmp = LLVMBuildGEP(builder, tmp, gep_idx, 3, "");
3502 tmp = LLVMBuildLoad(builder, tmp, "");
3503 prim.isnull = LLVMBuildICmp(builder, LLVMIntEQ, tmp,
3504 LLVMConstInt(ctx->ac.i8, 0, false), "");
3505
3506 for (unsigned i = 0; i < verts_per_prim; ++i) {
3507 prim.index[i] = LLVMBuildSub(builder, vertlive_scan.result_exclusive,
3508 LLVMConstInt(ctx->ac.i32, verts_per_prim - i - 1, false), "");
3509 prim.edgeflag[i] = ctx->ac.i1false;
3510 }
3511
3512 build_export_prim(ctx, &prim);
3513 }
3514 ac_build_endif(&ctx->ac, 5140);
3515
3516 /* Export position and parameter data */
3517 tmp = LLVMBuildICmp(builder, LLVMIntULT, tid, vertlive_scan.result_reduce, "");
3518 ac_build_ifcc(&ctx->ac, tmp, 5145);
3519 {
3520 struct radv_vs_output_info *outinfo = &ctx->shader_info->vs.outinfo;
3521 bool export_view_index = ctx->options->key.has_multiview_view_index;
3522 struct radv_shader_output_values *outputs;
3523 unsigned noutput = 0;
3524
3525 /* Allocate a temporary array for the output values. */
3526 unsigned num_outputs = util_bitcount64(ctx->output_mask) + export_view_index;
3527 outputs = calloc(num_outputs, sizeof(outputs[0]));
3528
3529 memset(outinfo->vs_output_param_offset, AC_EXP_PARAM_UNDEFINED,
3530 sizeof(outinfo->vs_output_param_offset));
3531 outinfo->pos_exports = 0;
3532
3533 tmp = ngg_gs_vertex_ptr(ctx, tid);
3534 LLVMValueRef gep_idx[3] = {
3535 ctx->ac.i32_0, /* implicit C-style array */
3536 ctx->ac.i32_1, /* second value of struct */
3537 ctx->ac.i32_1, /* stream 1: source data index */
3538 };
3539 tmp = LLVMBuildGEP(builder, tmp, gep_idx, 3, "");
3540 tmp = LLVMBuildLoad(builder, tmp, "");
3541 tmp = LLVMBuildZExt(builder, tmp, ctx->ac.i32, "");
3542 const LLVMValueRef vertexptr = ngg_gs_vertex_ptr(ctx, tmp);
3543
3544 if (ctx->output_mask & (1ull << VARYING_SLOT_PSIZ)) {
3545 outinfo->writes_pointsize = true;
3546 }
3547
3548 if (ctx->output_mask & (1ull << VARYING_SLOT_LAYER)) {
3549 outinfo->writes_layer = true;
3550 }
3551
3552 if (ctx->output_mask & (1ull << VARYING_SLOT_VIEWPORT)) {
3553 outinfo->writes_viewport_index = true;
3554 }
3555
3556 unsigned out_idx = 0;
3557 gep_idx[1] = ctx->ac.i32_0;
3558 for (unsigned i = 0; i < AC_LLVM_MAX_OUTPUTS; ++i) {
3559 if (!(ctx->output_mask & (1ull << i)))
3560 continue;
3561
3562 outputs[noutput].slot_name = i;
3563 outputs[noutput].slot_index = i == VARYING_SLOT_CLIP_DIST1;
3564
3565 outputs[noutput].usage_mask = ctx->shader_info->info.gs.output_usage_mask[i];
3566 int length = util_last_bit(outputs[noutput].usage_mask);
3567
3568 for (unsigned j = 0; j < length; j++, out_idx++) {
3569 gep_idx[2] = LLVMConstInt(ctx->ac.i32, out_idx, false);
3570 tmp = LLVMBuildGEP(builder, vertexptr, gep_idx, 3, "");
3571 tmp = LLVMBuildLoad(builder, tmp, "");
3572
3573 LLVMTypeRef type = LLVMGetAllocatedType(ctx->abi.outputs[ac_llvm_reg_index_soa(i, j)]);
3574 if (ac_get_type_size(type) == 2) {
3575 tmp = ac_to_integer(&ctx->ac, tmp);
3576 tmp = LLVMBuildTrunc(ctx->ac.builder, tmp, ctx->ac.i16, "");
3577 }
3578
3579 outputs[noutput].values[j] = ac_to_float(&ctx->ac, tmp);
3580 }
3581
3582 for (unsigned j = length; j < 4; j++)
3583 outputs[noutput].values[j] = LLVMGetUndef(ctx->ac.f32);
3584
3585 noutput++;
3586 }
3587
3588 /* Export ViewIndex. */
3589 if (export_view_index) {
3590 outinfo->writes_layer = true;
3591
3592 outputs[noutput].slot_name = VARYING_SLOT_LAYER;
3593 outputs[noutput].slot_index = 0;
3594 outputs[noutput].usage_mask = 0x1;
3595 outputs[noutput].values[0] = ac_to_float(&ctx->ac, ctx->abi.view_index);
3596 for (unsigned j = 1; j < 4; j++)
3597 outputs[noutput].values[j] = ctx->ac.f32_0;
3598 noutput++;
3599 }
3600
3601 radv_llvm_export_vs(ctx, outputs, noutput, outinfo,
3602 ctx->options->key.vs_common_out.export_clip_dists);
3603 FREE(outputs);
3604 }
3605 ac_build_endif(&ctx->ac, 5145);
3606 }
3607
3608 static void gfx10_ngg_gs_emit_vertex(struct radv_shader_context *ctx,
3609 unsigned stream,
3610 LLVMValueRef *addrs)
3611 {
3612 LLVMBuilderRef builder = ctx->ac.builder;
3613 LLVMValueRef tmp;
3614 const LLVMValueRef vertexidx =
3615 LLVMBuildLoad(builder, ctx->gs_next_vertex[stream], "");
3616
3617 /* If this thread has already emitted the declared maximum number of
3618 * vertices, skip the write: excessive vertex emissions are not
3619 * supposed to have any effect.
3620 */
3621 const LLVMValueRef can_emit =
3622 LLVMBuildICmp(builder, LLVMIntULT, vertexidx,
3623 LLVMConstInt(ctx->ac.i32, ctx->gs_max_out_vertices, false), "");
3624 ac_build_kill_if_false(&ctx->ac, can_emit);
3625
3626 tmp = LLVMBuildAdd(builder, vertexidx, ctx->ac.i32_1, "");
3627 tmp = LLVMBuildSelect(builder, can_emit, tmp, vertexidx, "");
3628 LLVMBuildStore(builder, tmp, ctx->gs_next_vertex[stream]);
3629
3630 const LLVMValueRef vertexptr =
3631 ngg_gs_emit_vertex_ptr(ctx, get_thread_id_in_tg(ctx), vertexidx);
3632 unsigned out_idx = 0;
3633 for (unsigned i = 0; i < AC_LLVM_MAX_OUTPUTS; ++i) {
3634 unsigned output_usage_mask =
3635 ctx->shader_info->info.gs.output_usage_mask[i];
3636 uint8_t output_stream =
3637 ctx->shader_info->info.gs.output_streams[i];
3638 LLVMValueRef *out_ptr = &addrs[i * 4];
3639 int length = util_last_bit(output_usage_mask);
3640
3641 if (!(ctx->output_mask & (1ull << i)) ||
3642 output_stream != stream)
3643 continue;
3644
3645 for (unsigned j = 0; j < length; j++, out_idx++) {
3646 if (!(output_usage_mask & (1 << j)))
3647 continue;
3648
3649 LLVMValueRef out_val = LLVMBuildLoad(ctx->ac.builder,
3650 out_ptr[j], "");
3651 LLVMValueRef gep_idx[3] = {
3652 ctx->ac.i32_0, /* implied C-style array */
3653 ctx->ac.i32_0, /* first entry of struct */
3654 LLVMConstInt(ctx->ac.i32, out_idx, false),
3655 };
3656 LLVMValueRef ptr = LLVMBuildGEP(builder, vertexptr, gep_idx, 3, "");
3657
3658 out_val = ac_to_integer(&ctx->ac, out_val);
3659 out_val = LLVMBuildZExtOrBitCast(ctx->ac.builder, out_val, ctx->ac.i32, "");
3660
3661 LLVMBuildStore(builder, out_val, ptr);
3662 }
3663 }
3664 assert(out_idx * 4 <= ctx->gsvs_vertex_size);
3665
3666 /* Determine and store whether this vertex completed a primitive. */
3667 const LLVMValueRef curverts = LLVMBuildLoad(builder, ctx->gs_curprim_verts[stream], "");
3668
3669 tmp = LLVMConstInt(ctx->ac.i32, si_conv_gl_prim_to_vertices(ctx->gs_output_prim) - 1, false);
3670 const LLVMValueRef iscompleteprim =
3671 LLVMBuildICmp(builder, LLVMIntUGE, curverts, tmp, "");
3672
3673 tmp = LLVMBuildAdd(builder, curverts, ctx->ac.i32_1, "");
3674 LLVMBuildStore(builder, tmp, ctx->gs_curprim_verts[stream]);
3675
3676 LLVMValueRef gep_idx[3] = {
3677 ctx->ac.i32_0, /* implied C-style array */
3678 ctx->ac.i32_1, /* second struct entry */
3679 LLVMConstInt(ctx->ac.i32, stream, false),
3680 };
3681 const LLVMValueRef primflagptr =
3682 LLVMBuildGEP(builder, vertexptr, gep_idx, 3, "");
3683
3684 tmp = LLVMBuildZExt(builder, iscompleteprim, ctx->ac.i8, "");
3685 LLVMBuildStore(builder, tmp, primflagptr);
3686
3687 tmp = LLVMBuildLoad(builder, ctx->gs_generated_prims[stream], "");
3688 tmp = LLVMBuildAdd(builder, tmp, LLVMBuildZExt(builder, iscompleteprim, ctx->ac.i32, ""), "");
3689 LLVMBuildStore(builder, tmp, ctx->gs_generated_prims[stream]);
3690 }
3691
3692 static void
3693 write_tess_factors(struct radv_shader_context *ctx)
3694 {
3695 unsigned stride, outer_comps, inner_comps;
3696 LLVMValueRef invocation_id = ac_unpack_param(&ctx->ac, ctx->abi.tcs_rel_ids, 8, 5);
3697 LLVMValueRef rel_patch_id = ac_unpack_param(&ctx->ac, ctx->abi.tcs_rel_ids, 0, 8);
3698 unsigned tess_inner_index = 0, tess_outer_index;
3699 LLVMValueRef lds_base, lds_inner = NULL, lds_outer, byteoffset, buffer;
3700 LLVMValueRef out[6], vec0, vec1, tf_base, inner[4], outer[4];
3701 int i;
3702 ac_emit_barrier(&ctx->ac, ctx->stage);
3703
3704 switch (ctx->options->key.tcs.primitive_mode) {
3705 case GL_ISOLINES:
3706 stride = 2;
3707 outer_comps = 2;
3708 inner_comps = 0;
3709 break;
3710 case GL_TRIANGLES:
3711 stride = 4;
3712 outer_comps = 3;
3713 inner_comps = 1;
3714 break;
3715 case GL_QUADS:
3716 stride = 6;
3717 outer_comps = 4;
3718 inner_comps = 2;
3719 break;
3720 default:
3721 return;
3722 }
3723
3724 ac_build_ifcc(&ctx->ac,
3725 LLVMBuildICmp(ctx->ac.builder, LLVMIntEQ,
3726 invocation_id, ctx->ac.i32_0, ""), 6503);
3727
3728 lds_base = get_tcs_out_current_patch_data_offset(ctx);
3729
3730 if (inner_comps) {
3731 tess_inner_index = shader_io_get_unique_index(VARYING_SLOT_TESS_LEVEL_INNER);
3732 lds_inner = LLVMBuildAdd(ctx->ac.builder, lds_base,
3733 LLVMConstInt(ctx->ac.i32, tess_inner_index * 4, false), "");
3734 }
3735
3736 tess_outer_index = shader_io_get_unique_index(VARYING_SLOT_TESS_LEVEL_OUTER);
3737 lds_outer = LLVMBuildAdd(ctx->ac.builder, lds_base,
3738 LLVMConstInt(ctx->ac.i32, tess_outer_index * 4, false), "");
3739
3740 for (i = 0; i < 4; i++) {
3741 inner[i] = LLVMGetUndef(ctx->ac.i32);
3742 outer[i] = LLVMGetUndef(ctx->ac.i32);
3743 }
3744
3745 // LINES reversal
3746 if (ctx->options->key.tcs.primitive_mode == GL_ISOLINES) {
3747 outer[0] = out[1] = ac_lds_load(&ctx->ac, lds_outer);
3748 lds_outer = LLVMBuildAdd(ctx->ac.builder, lds_outer,
3749 ctx->ac.i32_1, "");
3750 outer[1] = out[0] = ac_lds_load(&ctx->ac, lds_outer);
3751 } else {
3752 for (i = 0; i < outer_comps; i++) {
3753 outer[i] = out[i] =
3754 ac_lds_load(&ctx->ac, lds_outer);
3755 lds_outer = LLVMBuildAdd(ctx->ac.builder, lds_outer,
3756 ctx->ac.i32_1, "");
3757 }
3758 for (i = 0; i < inner_comps; i++) {
3759 inner[i] = out[outer_comps+i] =
3760 ac_lds_load(&ctx->ac, lds_inner);
3761 lds_inner = LLVMBuildAdd(ctx->ac.builder, lds_inner,
3762 ctx->ac.i32_1, "");
3763 }
3764 }
3765
3766 /* Convert the outputs to vectors for stores. */
3767 vec0 = ac_build_gather_values(&ctx->ac, out, MIN2(stride, 4));
3768 vec1 = NULL;
3769
3770 if (stride > 4)
3771 vec1 = ac_build_gather_values(&ctx->ac, out + 4, stride - 4);
3772
3773
3774 buffer = ctx->hs_ring_tess_factor;
3775 tf_base = ctx->tess_factor_offset;
3776 byteoffset = LLVMBuildMul(ctx->ac.builder, rel_patch_id,
3777 LLVMConstInt(ctx->ac.i32, 4 * stride, false), "");
3778 unsigned tf_offset = 0;
3779
3780 if (ctx->options->chip_class <= GFX8) {
3781 ac_build_ifcc(&ctx->ac,
3782 LLVMBuildICmp(ctx->ac.builder, LLVMIntEQ,
3783 rel_patch_id, ctx->ac.i32_0, ""), 6504);
3784
3785 /* Store the dynamic HS control word. */
3786 ac_build_buffer_store_dword(&ctx->ac, buffer,
3787 LLVMConstInt(ctx->ac.i32, 0x80000000, false),
3788 1, ctx->ac.i32_0, tf_base,
3789 0, ac_glc, false);
3790 tf_offset += 4;
3791
3792 ac_build_endif(&ctx->ac, 6504);
3793 }
3794
3795 /* Store the tessellation factors. */
3796 ac_build_buffer_store_dword(&ctx->ac, buffer, vec0,
3797 MIN2(stride, 4), byteoffset, tf_base,
3798 tf_offset, ac_glc, false);
3799 if (vec1)
3800 ac_build_buffer_store_dword(&ctx->ac, buffer, vec1,
3801 stride - 4, byteoffset, tf_base,
3802 16 + tf_offset, ac_glc, false);
3803
3804 //store to offchip for TES to read - only if TES reads them
3805 if (ctx->options->key.tcs.tes_reads_tess_factors) {
3806 LLVMValueRef inner_vec, outer_vec, tf_outer_offset;
3807 LLVMValueRef tf_inner_offset;
3808 unsigned param_outer, param_inner;
3809
3810 param_outer = shader_io_get_unique_index(VARYING_SLOT_TESS_LEVEL_OUTER);
3811 tf_outer_offset = get_tcs_tes_buffer_address(ctx, NULL,
3812 LLVMConstInt(ctx->ac.i32, param_outer, 0));
3813
3814 outer_vec = ac_build_gather_values(&ctx->ac, outer,
3815 util_next_power_of_two(outer_comps));
3816
3817 ac_build_buffer_store_dword(&ctx->ac, ctx->hs_ring_tess_offchip, outer_vec,
3818 outer_comps, tf_outer_offset,
3819 ctx->oc_lds, 0, ac_glc, false);
3820 if (inner_comps) {
3821 param_inner = shader_io_get_unique_index(VARYING_SLOT_TESS_LEVEL_INNER);
3822 tf_inner_offset = get_tcs_tes_buffer_address(ctx, NULL,
3823 LLVMConstInt(ctx->ac.i32, param_inner, 0));
3824
3825 inner_vec = inner_comps == 1 ? inner[0] :
3826 ac_build_gather_values(&ctx->ac, inner, inner_comps);
3827 ac_build_buffer_store_dword(&ctx->ac, ctx->hs_ring_tess_offchip, inner_vec,
3828 inner_comps, tf_inner_offset,
3829 ctx->oc_lds, 0, ac_glc, false);
3830 }
3831 }
3832
3833 ac_build_endif(&ctx->ac, 6503);
3834 }
3835
3836 static void
3837 handle_tcs_outputs_post(struct radv_shader_context *ctx)
3838 {
3839 write_tess_factors(ctx);
3840 }
3841
3842 static bool
3843 si_export_mrt_color(struct radv_shader_context *ctx,
3844 LLVMValueRef *color, unsigned index,
3845 struct ac_export_args *args)
3846 {
3847 /* Export */
3848 si_llvm_init_export_args(ctx, color, 0xf,
3849 V_008DFC_SQ_EXP_MRT + index, args);
3850 if (!args->enabled_channels)
3851 return false; /* unnecessary NULL export */
3852
3853 return true;
3854 }
3855
3856 static void
3857 radv_export_mrt_z(struct radv_shader_context *ctx,
3858 LLVMValueRef depth, LLVMValueRef stencil,
3859 LLVMValueRef samplemask)
3860 {
3861 struct ac_export_args args;
3862
3863 ac_export_mrt_z(&ctx->ac, depth, stencil, samplemask, &args);
3864
3865 ac_build_export(&ctx->ac, &args);
3866 }
3867
3868 static void
3869 handle_fs_outputs_post(struct radv_shader_context *ctx)
3870 {
3871 unsigned index = 0;
3872 LLVMValueRef depth = NULL, stencil = NULL, samplemask = NULL;
3873 struct ac_export_args color_args[8];
3874
3875 for (unsigned i = 0; i < AC_LLVM_MAX_OUTPUTS; ++i) {
3876 LLVMValueRef values[4];
3877
3878 if (!(ctx->output_mask & (1ull << i)))
3879 continue;
3880
3881 if (i < FRAG_RESULT_DATA0)
3882 continue;
3883
3884 for (unsigned j = 0; j < 4; j++)
3885 values[j] = ac_to_float(&ctx->ac,
3886 radv_load_output(ctx, i, j));
3887
3888 bool ret = si_export_mrt_color(ctx, values,
3889 i - FRAG_RESULT_DATA0,
3890 &color_args[index]);
3891 if (ret)
3892 index++;
3893 }
3894
3895 /* Process depth, stencil, samplemask. */
3896 if (ctx->shader_info->info.ps.writes_z) {
3897 depth = ac_to_float(&ctx->ac,
3898 radv_load_output(ctx, FRAG_RESULT_DEPTH, 0));
3899 }
3900 if (ctx->shader_info->info.ps.writes_stencil) {
3901 stencil = ac_to_float(&ctx->ac,
3902 radv_load_output(ctx, FRAG_RESULT_STENCIL, 0));
3903 }
3904 if (ctx->shader_info->info.ps.writes_sample_mask) {
3905 samplemask = ac_to_float(&ctx->ac,
3906 radv_load_output(ctx, FRAG_RESULT_SAMPLE_MASK, 0));
3907 }
3908
3909 /* Set the DONE bit on last non-null color export only if Z isn't
3910 * exported.
3911 */
3912 if (index > 0 &&
3913 !ctx->shader_info->info.ps.writes_z &&
3914 !ctx->shader_info->info.ps.writes_stencil &&
3915 !ctx->shader_info->info.ps.writes_sample_mask) {
3916 unsigned last = index - 1;
3917
3918 color_args[last].valid_mask = 1; /* whether the EXEC mask is valid */
3919 color_args[last].done = 1; /* DONE bit */
3920 }
3921
3922 /* Export PS outputs. */
3923 for (unsigned i = 0; i < index; i++)
3924 ac_build_export(&ctx->ac, &color_args[i]);
3925
3926 if (depth || stencil || samplemask)
3927 radv_export_mrt_z(ctx, depth, stencil, samplemask);
3928 else if (!index)
3929 ac_build_export_null(&ctx->ac);
3930 }
3931
3932 static void
3933 emit_gs_epilogue(struct radv_shader_context *ctx)
3934 {
3935 if (ctx->options->key.vs_common_out.as_ngg) {
3936 gfx10_ngg_gs_emit_epilogue_1(ctx);
3937 return;
3938 }
3939
3940 if (ctx->ac.chip_class >= GFX10)
3941 LLVMBuildFence(ctx->ac.builder, LLVMAtomicOrderingRelease, false, "");
3942
3943 ac_build_sendmsg(&ctx->ac, AC_SENDMSG_GS_OP_NOP | AC_SENDMSG_GS_DONE, ctx->gs_wave_id);
3944 }
3945
3946 static void
3947 handle_shader_outputs_post(struct ac_shader_abi *abi, unsigned max_outputs,
3948 LLVMValueRef *addrs)
3949 {
3950 struct radv_shader_context *ctx = radv_shader_context_from_abi(abi);
3951
3952 switch (ctx->stage) {
3953 case MESA_SHADER_VERTEX:
3954 if (ctx->options->key.vs_common_out.as_ls)
3955 handle_ls_outputs_post(ctx);
3956 else if (ctx->options->key.vs_common_out.as_es)
3957 handle_es_outputs_post(ctx, &ctx->shader_info->vs.es_info);
3958 else if (ctx->options->key.vs_common_out.as_ngg)
3959 break; /* handled outside of the shader body */
3960 else
3961 handle_vs_outputs_post(ctx, ctx->options->key.vs_common_out.export_prim_id,
3962 ctx->options->key.vs_common_out.export_clip_dists,
3963 &ctx->shader_info->vs.outinfo);
3964 break;
3965 case MESA_SHADER_FRAGMENT:
3966 handle_fs_outputs_post(ctx);
3967 break;
3968 case MESA_SHADER_GEOMETRY:
3969 emit_gs_epilogue(ctx);
3970 break;
3971 case MESA_SHADER_TESS_CTRL:
3972 handle_tcs_outputs_post(ctx);
3973 break;
3974 case MESA_SHADER_TESS_EVAL:
3975 if (ctx->options->key.vs_common_out.as_es)
3976 handle_es_outputs_post(ctx, &ctx->shader_info->tes.es_info);
3977 else if (ctx->options->key.vs_common_out.as_ngg)
3978 break; /* handled outside of the shader body */
3979 else
3980 handle_vs_outputs_post(ctx, ctx->options->key.vs_common_out.export_prim_id,
3981 ctx->options->key.vs_common_out.export_clip_dists,
3982 &ctx->shader_info->tes.outinfo);
3983 break;
3984 default:
3985 break;
3986 }
3987 }
3988
3989 static void ac_llvm_finalize_module(struct radv_shader_context *ctx,
3990 LLVMPassManagerRef passmgr,
3991 const struct radv_nir_compiler_options *options)
3992 {
3993 LLVMRunPassManager(passmgr, ctx->ac.module);
3994 LLVMDisposeBuilder(ctx->ac.builder);
3995
3996 ac_llvm_context_dispose(&ctx->ac);
3997 }
3998
3999 static void
4000 ac_nir_eliminate_const_vs_outputs(struct radv_shader_context *ctx)
4001 {
4002 struct radv_vs_output_info *outinfo;
4003
4004 switch (ctx->stage) {
4005 case MESA_SHADER_FRAGMENT:
4006 case MESA_SHADER_COMPUTE:
4007 case MESA_SHADER_TESS_CTRL:
4008 case MESA_SHADER_GEOMETRY:
4009 return;
4010 case MESA_SHADER_VERTEX:
4011 if (ctx->options->key.vs_common_out.as_ls ||
4012 ctx->options->key.vs_common_out.as_es)
4013 return;
4014 outinfo = &ctx->shader_info->vs.outinfo;
4015 break;
4016 case MESA_SHADER_TESS_EVAL:
4017 if (ctx->options->key.vs_common_out.as_es)
4018 return;
4019 outinfo = &ctx->shader_info->tes.outinfo;
4020 break;
4021 default:
4022 unreachable("Unhandled shader type");
4023 }
4024
4025 ac_optimize_vs_outputs(&ctx->ac,
4026 ctx->main_function,
4027 outinfo->vs_output_param_offset,
4028 VARYING_SLOT_MAX,
4029 &outinfo->param_exports);
4030 }
4031
4032 static void
4033 ac_setup_rings(struct radv_shader_context *ctx)
4034 {
4035 if (ctx->options->chip_class <= GFX8 &&
4036 (ctx->stage == MESA_SHADER_GEOMETRY ||
4037 ctx->options->key.vs_common_out.as_es || ctx->options->key.vs_common_out.as_es)) {
4038 unsigned ring = ctx->stage == MESA_SHADER_GEOMETRY ? RING_ESGS_GS
4039 : RING_ESGS_VS;
4040 LLVMValueRef offset = LLVMConstInt(ctx->ac.i32, ring, false);
4041
4042 ctx->esgs_ring = ac_build_load_to_sgpr(&ctx->ac,
4043 ctx->ring_offsets,
4044 offset);
4045 }
4046
4047 if (ctx->is_gs_copy_shader) {
4048 ctx->gsvs_ring[0] =
4049 ac_build_load_to_sgpr(&ctx->ac, ctx->ring_offsets,
4050 LLVMConstInt(ctx->ac.i32,
4051 RING_GSVS_VS, false));
4052 }
4053
4054 if (ctx->stage == MESA_SHADER_GEOMETRY) {
4055 /* The conceptual layout of the GSVS ring is
4056 * v0c0 .. vLv0 v0c1 .. vLc1 ..
4057 * but the real memory layout is swizzled across
4058 * threads:
4059 * t0v0c0 .. t15v0c0 t0v1c0 .. t15v1c0 ... t15vLcL
4060 * t16v0c0 ..
4061 * Override the buffer descriptor accordingly.
4062 */
4063 LLVMTypeRef v2i64 = LLVMVectorType(ctx->ac.i64, 2);
4064 uint64_t stream_offset = 0;
4065 unsigned num_records = ctx->ac.wave_size;
4066 LLVMValueRef base_ring;
4067
4068 base_ring =
4069 ac_build_load_to_sgpr(&ctx->ac, ctx->ring_offsets,
4070 LLVMConstInt(ctx->ac.i32,
4071 RING_GSVS_GS, false));
4072
4073 for (unsigned stream = 0; stream < 4; stream++) {
4074 unsigned num_components, stride;
4075 LLVMValueRef ring, tmp;
4076
4077 num_components =
4078 ctx->shader_info->info.gs.num_stream_output_components[stream];
4079
4080 if (!num_components)
4081 continue;
4082
4083 stride = 4 * num_components * ctx->gs_max_out_vertices;
4084
4085 /* Limit on the stride field for <= GFX7. */
4086 assert(stride < (1 << 14));
4087
4088 ring = LLVMBuildBitCast(ctx->ac.builder,
4089 base_ring, v2i64, "");
4090 tmp = LLVMBuildExtractElement(ctx->ac.builder,
4091 ring, ctx->ac.i32_0, "");
4092 tmp = LLVMBuildAdd(ctx->ac.builder, tmp,
4093 LLVMConstInt(ctx->ac.i64,
4094 stream_offset, 0), "");
4095 ring = LLVMBuildInsertElement(ctx->ac.builder,
4096 ring, tmp, ctx->ac.i32_0, "");
4097
4098 stream_offset += stride * ctx->ac.wave_size;
4099
4100 ring = LLVMBuildBitCast(ctx->ac.builder, ring,
4101 ctx->ac.v4i32, "");
4102
4103 tmp = LLVMBuildExtractElement(ctx->ac.builder, ring,
4104 ctx->ac.i32_1, "");
4105 tmp = LLVMBuildOr(ctx->ac.builder, tmp,
4106 LLVMConstInt(ctx->ac.i32,
4107 S_008F04_STRIDE(stride), false), "");
4108 ring = LLVMBuildInsertElement(ctx->ac.builder, ring, tmp,
4109 ctx->ac.i32_1, "");
4110
4111 ring = LLVMBuildInsertElement(ctx->ac.builder, ring,
4112 LLVMConstInt(ctx->ac.i32,
4113 num_records, false),
4114 LLVMConstInt(ctx->ac.i32, 2, false), "");
4115
4116 ctx->gsvs_ring[stream] = ring;
4117 }
4118 }
4119
4120 if (ctx->stage == MESA_SHADER_TESS_CTRL ||
4121 ctx->stage == MESA_SHADER_TESS_EVAL) {
4122 ctx->hs_ring_tess_offchip = ac_build_load_to_sgpr(&ctx->ac, ctx->ring_offsets, LLVMConstInt(ctx->ac.i32, RING_HS_TESS_OFFCHIP, false));
4123 ctx->hs_ring_tess_factor = ac_build_load_to_sgpr(&ctx->ac, ctx->ring_offsets, LLVMConstInt(ctx->ac.i32, RING_HS_TESS_FACTOR, false));
4124 }
4125 }
4126
4127 unsigned
4128 radv_nir_get_max_workgroup_size(enum chip_class chip_class,
4129 gl_shader_stage stage,
4130 const struct nir_shader *nir)
4131 {
4132 const unsigned backup_sizes[] = {chip_class >= GFX9 ? 128 : 64, 1, 1};
4133 return radv_get_max_workgroup_size(chip_class, stage, nir ? nir->info.cs.local_size : backup_sizes);
4134 }
4135
4136 /* Fixup the HW not emitting the TCS regs if there are no HS threads. */
4137 static void ac_nir_fixup_ls_hs_input_vgprs(struct radv_shader_context *ctx)
4138 {
4139 LLVMValueRef count = ac_unpack_param(&ctx->ac, ctx->merged_wave_info, 8, 8);
4140 LLVMValueRef hs_empty = LLVMBuildICmp(ctx->ac.builder, LLVMIntEQ, count,
4141 ctx->ac.i32_0, "");
4142 ctx->abi.instance_id = LLVMBuildSelect(ctx->ac.builder, hs_empty, ctx->rel_auto_id, ctx->abi.instance_id, "");
4143 ctx->rel_auto_id = LLVMBuildSelect(ctx->ac.builder, hs_empty, ctx->abi.tcs_rel_ids, ctx->rel_auto_id, "");
4144 ctx->abi.vertex_id = LLVMBuildSelect(ctx->ac.builder, hs_empty, ctx->abi.tcs_patch_id, ctx->abi.vertex_id, "");
4145 }
4146
4147 static void prepare_gs_input_vgprs(struct radv_shader_context *ctx)
4148 {
4149 for(int i = 5; i >= 0; --i) {
4150 ctx->gs_vtx_offset[i] = ac_unpack_param(&ctx->ac, ctx->gs_vtx_offset[i & ~1],
4151 (i & 1) * 16, 16);
4152 }
4153
4154 ctx->gs_wave_id = ac_unpack_param(&ctx->ac, ctx->merged_wave_info, 16, 8);
4155 }
4156
4157 /* Ensure that the esgs ring is declared.
4158 *
4159 * We declare it with 64KB alignment as a hint that the
4160 * pointer value will always be 0.
4161 */
4162 static void declare_esgs_ring(struct radv_shader_context *ctx)
4163 {
4164 if (ctx->esgs_ring)
4165 return;
4166
4167 assert(!LLVMGetNamedGlobal(ctx->ac.module, "esgs_ring"));
4168
4169 ctx->esgs_ring = LLVMAddGlobalInAddressSpace(
4170 ctx->ac.module, LLVMArrayType(ctx->ac.i32, 0),
4171 "esgs_ring",
4172 AC_ADDR_SPACE_LDS);
4173 LLVMSetLinkage(ctx->esgs_ring, LLVMExternalLinkage);
4174 LLVMSetAlignment(ctx->esgs_ring, 64 * 1024);
4175 }
4176
4177 static
4178 LLVMModuleRef ac_translate_nir_to_llvm(struct ac_llvm_compiler *ac_llvm,
4179 struct nir_shader *const *shaders,
4180 int shader_count,
4181 struct radv_shader_variant_info *shader_info,
4182 const struct radv_nir_compiler_options *options)
4183 {
4184 struct radv_shader_context ctx = {0};
4185 unsigned i;
4186 ctx.options = options;
4187 ctx.shader_info = shader_info;
4188
4189 enum ac_float_mode float_mode =
4190 options->unsafe_math ? AC_FLOAT_MODE_UNSAFE_FP_MATH :
4191 AC_FLOAT_MODE_DEFAULT;
4192
4193 ac_llvm_context_init(&ctx.ac, ac_llvm, options->chip_class,
4194 options->family, float_mode, options->wave_size,
4195 options->wave_size);
4196 ctx.context = ctx.ac.context;
4197
4198 radv_nir_shader_info_init(&shader_info->info);
4199
4200 for(int i = 0; i < shader_count; ++i)
4201 radv_nir_shader_info_pass(shaders[i], options, &shader_info->info);
4202
4203 for (i = 0; i < MAX_SETS; i++)
4204 shader_info->user_sgprs_locs.descriptor_sets[i].sgpr_idx = -1;
4205 for (i = 0; i < AC_UD_MAX_UD; i++)
4206 shader_info->user_sgprs_locs.shader_data[i].sgpr_idx = -1;
4207
4208 ctx.max_workgroup_size = 0;
4209 for (int i = 0; i < shader_count; ++i) {
4210 ctx.max_workgroup_size = MAX2(ctx.max_workgroup_size,
4211 radv_nir_get_max_workgroup_size(ctx.options->chip_class,
4212 shaders[i]->info.stage,
4213 shaders[i]));
4214 }
4215
4216 if (ctx.ac.chip_class >= GFX10) {
4217 if (is_pre_gs_stage(shaders[0]->info.stage) &&
4218 options->key.vs_common_out.as_ngg) {
4219 ctx.max_workgroup_size = 128;
4220 }
4221 }
4222
4223 create_function(&ctx, shaders[shader_count - 1]->info.stage, shader_count >= 2,
4224 shader_count >= 2 ? shaders[shader_count - 2]->info.stage : MESA_SHADER_VERTEX);
4225
4226 ctx.abi.inputs = &ctx.inputs[0];
4227 ctx.abi.emit_outputs = handle_shader_outputs_post;
4228 ctx.abi.emit_vertex = visit_emit_vertex;
4229 ctx.abi.load_ubo = radv_load_ubo;
4230 ctx.abi.load_ssbo = radv_load_ssbo;
4231 ctx.abi.load_sampler_desc = radv_get_sampler_desc;
4232 ctx.abi.load_resource = radv_load_resource;
4233 ctx.abi.clamp_shadow_reference = false;
4234 ctx.abi.robust_buffer_access = options->robust_buffer_access;
4235
4236 bool is_ngg = is_pre_gs_stage(shaders[0]->info.stage) && ctx.options->key.vs_common_out.as_ngg;
4237 if (shader_count >= 2 || is_ngg)
4238 ac_init_exec_full_mask(&ctx.ac);
4239
4240 if (options->has_ls_vgpr_init_bug &&
4241 shaders[shader_count - 1]->info.stage == MESA_SHADER_TESS_CTRL)
4242 ac_nir_fixup_ls_hs_input_vgprs(&ctx);
4243
4244 for(int i = 0; i < shader_count; ++i) {
4245 ctx.stage = shaders[i]->info.stage;
4246 ctx.output_mask = 0;
4247
4248 if (shaders[i]->info.stage == MESA_SHADER_GEOMETRY) {
4249 for (int i = 0; i < 4; i++) {
4250 ctx.gs_next_vertex[i] =
4251 ac_build_alloca(&ctx.ac, ctx.ac.i32, "");
4252 }
4253 if (ctx.options->key.vs_common_out.as_ngg) {
4254 for (unsigned i = 0; i < 4; ++i) {
4255 ctx.gs_curprim_verts[i] =
4256 ac_build_alloca(&ctx.ac, ctx.ac.i32, "");
4257 ctx.gs_generated_prims[i] =
4258 ac_build_alloca(&ctx.ac, ctx.ac.i32, "");
4259 }
4260
4261 /* TODO: streamout */
4262
4263 LLVMTypeRef ai32 = LLVMArrayType(ctx.ac.i32, 8);
4264 ctx.gs_ngg_scratch =
4265 LLVMAddGlobalInAddressSpace(ctx.ac.module,
4266 ai32, "ngg_scratch", AC_ADDR_SPACE_LDS);
4267 LLVMSetInitializer(ctx.gs_ngg_scratch, LLVMGetUndef(ai32));
4268 LLVMSetAlignment(ctx.gs_ngg_scratch, 4);
4269
4270 ctx.gs_ngg_emit = LLVMBuildIntToPtr(ctx.ac.builder, ctx.ac.i32_0,
4271 LLVMPointerType(LLVMArrayType(ctx.ac.i32, 0), AC_ADDR_SPACE_LDS),
4272 "ngg_emit");
4273 }
4274
4275 ctx.gs_max_out_vertices = shaders[i]->info.gs.vertices_out;
4276 ctx.gs_output_prim = shaders[i]->info.gs.output_primitive;
4277 ctx.abi.load_inputs = load_gs_input;
4278 ctx.abi.emit_primitive = visit_end_primitive;
4279 } else if (shaders[i]->info.stage == MESA_SHADER_TESS_CTRL) {
4280 ctx.tcs_outputs_read = shaders[i]->info.outputs_read;
4281 ctx.tcs_patch_outputs_read = shaders[i]->info.patch_outputs_read;
4282 ctx.abi.load_tess_varyings = load_tcs_varyings;
4283 ctx.abi.load_patch_vertices_in = load_patch_vertices_in;
4284 ctx.abi.store_tcs_outputs = store_tcs_output;
4285 ctx.tcs_vertices_per_patch = shaders[i]->info.tess.tcs_vertices_out;
4286 if (shader_count == 1)
4287 ctx.tcs_num_inputs = ctx.options->key.tcs.num_inputs;
4288 else
4289 ctx.tcs_num_inputs = util_last_bit64(shader_info->info.vs.ls_outputs_written);
4290 ctx.tcs_num_patches = get_tcs_num_patches(&ctx);
4291 } else if (shaders[i]->info.stage == MESA_SHADER_TESS_EVAL) {
4292 ctx.tes_primitive_mode = shaders[i]->info.tess.primitive_mode;
4293 ctx.abi.load_tess_varyings = load_tes_input;
4294 ctx.abi.load_tess_coord = load_tess_coord;
4295 ctx.abi.load_patch_vertices_in = load_patch_vertices_in;
4296 ctx.tcs_vertices_per_patch = shaders[i]->info.tess.tcs_vertices_out;
4297 ctx.tcs_num_patches = ctx.options->key.tes.num_patches;
4298 } else if (shaders[i]->info.stage == MESA_SHADER_VERTEX) {
4299 ctx.abi.load_base_vertex = radv_load_base_vertex;
4300 } else if (shaders[i]->info.stage == MESA_SHADER_FRAGMENT) {
4301 ctx.abi.load_sample_position = load_sample_position;
4302 ctx.abi.load_sample_mask_in = load_sample_mask_in;
4303 ctx.abi.emit_kill = radv_emit_kill;
4304 }
4305
4306 if (shaders[i]->info.stage == MESA_SHADER_VERTEX &&
4307 ctx.options->key.vs_common_out.as_ngg &&
4308 ctx.options->key.vs_common_out.export_prim_id) {
4309 declare_esgs_ring(&ctx);
4310 }
4311
4312 bool nested_barrier = false;
4313
4314 if (i) {
4315 if (shaders[i]->info.stage == MESA_SHADER_GEOMETRY &&
4316 ctx.options->key.vs_common_out.as_ngg) {
4317 gfx10_ngg_gs_emit_prologue(&ctx);
4318 nested_barrier = false;
4319 } else {
4320 nested_barrier = true;
4321 }
4322 }
4323
4324 if (nested_barrier) {
4325 /* Execute a barrier before the second shader in
4326 * a merged shader.
4327 *
4328 * Execute the barrier inside the conditional block,
4329 * so that empty waves can jump directly to s_endpgm,
4330 * which will also signal the barrier.
4331 *
4332 * This is possible in gfx9, because an empty wave
4333 * for the second shader does not participate in
4334 * the epilogue. With NGG, empty waves may still
4335 * be required to export data (e.g. GS output vertices),
4336 * so we cannot let them exit early.
4337 *
4338 * If the shader is TCS and the TCS epilog is present
4339 * and contains a barrier, it will wait there and then
4340 * reach s_endpgm.
4341 */
4342 ac_emit_barrier(&ctx.ac, ctx.stage);
4343 }
4344
4345 nir_foreach_variable(variable, &shaders[i]->outputs)
4346 scan_shader_output_decl(&ctx, variable, shaders[i], shaders[i]->info.stage);
4347
4348 if (shaders[i]->info.stage == MESA_SHADER_GEOMETRY) {
4349 unsigned addclip = shaders[i]->info.clip_distance_array_size +
4350 shaders[i]->info.cull_distance_array_size > 4;
4351 ctx.gsvs_vertex_size = (util_bitcount64(ctx.output_mask) + addclip) * 16;
4352 ctx.max_gsvs_emit_size = ctx.gsvs_vertex_size *
4353 shaders[i]->info.gs.vertices_out;
4354 }
4355
4356 ac_setup_rings(&ctx);
4357
4358 LLVMBasicBlockRef merge_block;
4359 if (shader_count >= 2 || is_ngg) {
4360 LLVMValueRef fn = LLVMGetBasicBlockParent(LLVMGetInsertBlock(ctx.ac.builder));
4361 LLVMBasicBlockRef then_block = LLVMAppendBasicBlockInContext(ctx.ac.context, fn, "");
4362 merge_block = LLVMAppendBasicBlockInContext(ctx.ac.context, fn, "");
4363
4364 LLVMValueRef count = ac_unpack_param(&ctx.ac, ctx.merged_wave_info, 8 * i, 8);
4365 LLVMValueRef thread_id = ac_get_thread_id(&ctx.ac);
4366 LLVMValueRef cond = LLVMBuildICmp(ctx.ac.builder, LLVMIntULT,
4367 thread_id, count, "");
4368 LLVMBuildCondBr(ctx.ac.builder, cond, then_block, merge_block);
4369
4370 LLVMPositionBuilderAtEnd(ctx.ac.builder, then_block);
4371 }
4372
4373 if (shaders[i]->info.stage == MESA_SHADER_FRAGMENT)
4374 prepare_interp_optimize(&ctx, shaders[i]);
4375 else if(shaders[i]->info.stage == MESA_SHADER_VERTEX)
4376 handle_vs_inputs(&ctx, shaders[i]);
4377 else if(shader_count >= 2 && shaders[i]->info.stage == MESA_SHADER_GEOMETRY)
4378 prepare_gs_input_vgprs(&ctx);
4379
4380 ac_nir_translate(&ctx.ac, &ctx.abi, shaders[i]);
4381
4382 if (shader_count >= 2 || is_ngg) {
4383 LLVMBuildBr(ctx.ac.builder, merge_block);
4384 LLVMPositionBuilderAtEnd(ctx.ac.builder, merge_block);
4385 }
4386
4387 /* This needs to be outside the if wrapping the shader body, as sometimes
4388 * the HW generates waves with 0 es/vs threads. */
4389 if (is_pre_gs_stage(shaders[i]->info.stage) &&
4390 ctx.options->key.vs_common_out.as_ngg &&
4391 i == shader_count - 1) {
4392 handle_ngg_outputs_post(&ctx);
4393 } else if (shaders[i]->info.stage == MESA_SHADER_GEOMETRY &&
4394 ctx.options->key.vs_common_out.as_ngg) {
4395 gfx10_ngg_gs_emit_epilogue_2(&ctx);
4396 }
4397
4398 if (shaders[i]->info.stage == MESA_SHADER_GEOMETRY) {
4399 shader_info->gs.gsvs_vertex_size = ctx.gsvs_vertex_size;
4400 shader_info->gs.max_gsvs_emit_size = ctx.max_gsvs_emit_size;
4401 } else if (shaders[i]->info.stage == MESA_SHADER_TESS_CTRL) {
4402 shader_info->tcs.num_patches = ctx.tcs_num_patches;
4403 shader_info->tcs.lds_size = calculate_tess_lds_size(&ctx);
4404 }
4405 }
4406
4407 LLVMBuildRetVoid(ctx.ac.builder);
4408
4409 if (options->dump_preoptir) {
4410 fprintf(stderr, "%s LLVM IR:\n\n",
4411 radv_get_shader_name(shader_info,
4412 shaders[shader_count - 1]->info.stage));
4413 ac_dump_module(ctx.ac.module);
4414 fprintf(stderr, "\n");
4415 }
4416
4417 ac_llvm_finalize_module(&ctx, ac_llvm->passmgr, options);
4418
4419 if (shader_count == 1)
4420 ac_nir_eliminate_const_vs_outputs(&ctx);
4421
4422 if (options->dump_shader) {
4423 ctx.shader_info->private_mem_vgprs =
4424 ac_count_scratch_private_memory(ctx.main_function);
4425 }
4426
4427 return ctx.ac.module;
4428 }
4429
4430 static void ac_diagnostic_handler(LLVMDiagnosticInfoRef di, void *context)
4431 {
4432 unsigned *retval = (unsigned *)context;
4433 LLVMDiagnosticSeverity severity = LLVMGetDiagInfoSeverity(di);
4434 char *description = LLVMGetDiagInfoDescription(di);
4435
4436 if (severity == LLVMDSError) {
4437 *retval = 1;
4438 fprintf(stderr, "LLVM triggered Diagnostic Handler: %s\n",
4439 description);
4440 }
4441
4442 LLVMDisposeMessage(description);
4443 }
4444
4445 static unsigned radv_llvm_compile(LLVMModuleRef M,
4446 char **pelf_buffer, size_t *pelf_size,
4447 struct ac_llvm_compiler *ac_llvm)
4448 {
4449 unsigned retval = 0;
4450 LLVMContextRef llvm_ctx;
4451
4452 /* Setup Diagnostic Handler*/
4453 llvm_ctx = LLVMGetModuleContext(M);
4454
4455 LLVMContextSetDiagnosticHandler(llvm_ctx, ac_diagnostic_handler,
4456 &retval);
4457
4458 /* Compile IR*/
4459 if (!radv_compile_to_elf(ac_llvm, M, pelf_buffer, pelf_size))
4460 retval = 1;
4461 return retval;
4462 }
4463
4464 static void ac_compile_llvm_module(struct ac_llvm_compiler *ac_llvm,
4465 LLVMModuleRef llvm_module,
4466 struct radv_shader_binary **rbinary,
4467 gl_shader_stage stage,
4468 const char *name,
4469 const struct radv_nir_compiler_options *options)
4470 {
4471 char *elf_buffer = NULL;
4472 size_t elf_size = 0;
4473 char *llvm_ir_string = NULL;
4474
4475 if (options->dump_shader) {
4476 fprintf(stderr, "%s LLVM IR:\n\n", name);
4477 ac_dump_module(llvm_module);
4478 fprintf(stderr, "\n");
4479 }
4480
4481 if (options->record_llvm_ir) {
4482 char *llvm_ir = LLVMPrintModuleToString(llvm_module);
4483 llvm_ir_string = strdup(llvm_ir);
4484 LLVMDisposeMessage(llvm_ir);
4485 }
4486
4487 int v = radv_llvm_compile(llvm_module, &elf_buffer, &elf_size, ac_llvm);
4488 if (v) {
4489 fprintf(stderr, "compile failed\n");
4490 }
4491
4492 LLVMContextRef ctx = LLVMGetModuleContext(llvm_module);
4493 LLVMDisposeModule(llvm_module);
4494 LLVMContextDispose(ctx);
4495
4496 size_t llvm_ir_size = llvm_ir_string ? strlen(llvm_ir_string) : 0;
4497 size_t alloc_size = sizeof(struct radv_shader_binary_rtld) + elf_size + llvm_ir_size + 1;
4498 struct radv_shader_binary_rtld *rbin = calloc(1, alloc_size);
4499 memcpy(rbin->data, elf_buffer, elf_size);
4500 if (llvm_ir_string)
4501 memcpy(rbin->data + elf_size, llvm_ir_string, llvm_ir_size + 1);
4502
4503 rbin->base.type = RADV_BINARY_TYPE_RTLD;
4504 rbin->base.stage = stage;
4505 rbin->base.total_size = alloc_size;
4506 rbin->elf_size = elf_size;
4507 rbin->llvm_ir_size = llvm_ir_size;
4508 *rbinary = &rbin->base;
4509
4510 free(llvm_ir_string);
4511 free(elf_buffer);
4512 }
4513
4514 static void
4515 ac_fill_shader_info(struct radv_shader_variant_info *shader_info, struct nir_shader *nir, const struct radv_nir_compiler_options *options)
4516 {
4517 switch (nir->info.stage) {
4518 case MESA_SHADER_COMPUTE:
4519 for (int i = 0; i < 3; ++i)
4520 shader_info->cs.block_size[i] = nir->info.cs.local_size[i];
4521 break;
4522 case MESA_SHADER_FRAGMENT:
4523 shader_info->fs.can_discard = nir->info.fs.uses_discard;
4524 shader_info->fs.early_fragment_test = nir->info.fs.early_fragment_tests;
4525 shader_info->fs.post_depth_coverage = nir->info.fs.post_depth_coverage;
4526 break;
4527 case MESA_SHADER_GEOMETRY:
4528 shader_info->gs.vertices_in = nir->info.gs.vertices_in;
4529 shader_info->gs.vertices_out = nir->info.gs.vertices_out;
4530 shader_info->gs.output_prim = nir->info.gs.output_primitive;
4531 shader_info->gs.invocations = nir->info.gs.invocations;
4532 break;
4533 case MESA_SHADER_TESS_EVAL:
4534 shader_info->tes.primitive_mode = nir->info.tess.primitive_mode;
4535 shader_info->tes.spacing = nir->info.tess.spacing;
4536 shader_info->tes.ccw = nir->info.tess.ccw;
4537 shader_info->tes.point_mode = nir->info.tess.point_mode;
4538 shader_info->tes.as_es = options->key.vs_common_out.as_es;
4539 shader_info->tes.export_prim_id = options->key.vs_common_out.export_prim_id;
4540 shader_info->is_ngg = options->key.vs_common_out.as_ngg;
4541 break;
4542 case MESA_SHADER_TESS_CTRL:
4543 shader_info->tcs.tcs_vertices_out = nir->info.tess.tcs_vertices_out;
4544 break;
4545 case MESA_SHADER_VERTEX:
4546 shader_info->vs.as_es = options->key.vs_common_out.as_es;
4547 shader_info->vs.as_ls = options->key.vs_common_out.as_ls;
4548 shader_info->vs.export_prim_id = options->key.vs_common_out.export_prim_id;
4549 shader_info->is_ngg = options->key.vs_common_out.as_ngg;
4550 break;
4551 default:
4552 break;
4553 }
4554 }
4555
4556 void
4557 radv_compile_nir_shader(struct ac_llvm_compiler *ac_llvm,
4558 struct radv_shader_binary **rbinary,
4559 struct radv_shader_variant_info *shader_info,
4560 struct nir_shader *const *nir,
4561 int nir_count,
4562 const struct radv_nir_compiler_options *options)
4563 {
4564
4565 LLVMModuleRef llvm_module;
4566
4567 llvm_module = ac_translate_nir_to_llvm(ac_llvm, nir, nir_count, shader_info,
4568 options);
4569
4570 ac_compile_llvm_module(ac_llvm, llvm_module, rbinary,
4571 nir[nir_count - 1]->info.stage,
4572 radv_get_shader_name(shader_info,
4573 nir[nir_count - 1]->info.stage),
4574 options);
4575
4576 for (int i = 0; i < nir_count; ++i)
4577 ac_fill_shader_info(shader_info, nir[i], options);
4578
4579 /* Determine the ES type (VS or TES) for the GS on GFX9. */
4580 if (options->chip_class >= GFX9) {
4581 if (nir_count == 2 &&
4582 nir[1]->info.stage == MESA_SHADER_GEOMETRY) {
4583 shader_info->gs.es_type = nir[0]->info.stage;
4584 }
4585 }
4586 shader_info->info.wave_size = options->wave_size;
4587 }
4588
4589 static void
4590 ac_gs_copy_shader_emit(struct radv_shader_context *ctx)
4591 {
4592 LLVMValueRef vtx_offset =
4593 LLVMBuildMul(ctx->ac.builder, ctx->abi.vertex_id,
4594 LLVMConstInt(ctx->ac.i32, 4, false), "");
4595 LLVMValueRef stream_id;
4596
4597 /* Fetch the vertex stream ID. */
4598 if (ctx->shader_info->info.so.num_outputs) {
4599 stream_id =
4600 ac_unpack_param(&ctx->ac, ctx->streamout_config, 24, 2);
4601 } else {
4602 stream_id = ctx->ac.i32_0;
4603 }
4604
4605 LLVMBasicBlockRef end_bb;
4606 LLVMValueRef switch_inst;
4607
4608 end_bb = LLVMAppendBasicBlockInContext(ctx->ac.context,
4609 ctx->main_function, "end");
4610 switch_inst = LLVMBuildSwitch(ctx->ac.builder, stream_id, end_bb, 4);
4611
4612 for (unsigned stream = 0; stream < 4; stream++) {
4613 unsigned num_components =
4614 ctx->shader_info->info.gs.num_stream_output_components[stream];
4615 LLVMBasicBlockRef bb;
4616 unsigned offset;
4617
4618 if (!num_components)
4619 continue;
4620
4621 if (stream > 0 && !ctx->shader_info->info.so.num_outputs)
4622 continue;
4623
4624 bb = LLVMInsertBasicBlockInContext(ctx->ac.context, end_bb, "out");
4625 LLVMAddCase(switch_inst, LLVMConstInt(ctx->ac.i32, stream, 0), bb);
4626 LLVMPositionBuilderAtEnd(ctx->ac.builder, bb);
4627
4628 offset = 0;
4629 for (unsigned i = 0; i < AC_LLVM_MAX_OUTPUTS; ++i) {
4630 unsigned output_usage_mask =
4631 ctx->shader_info->info.gs.output_usage_mask[i];
4632 unsigned output_stream =
4633 ctx->shader_info->info.gs.output_streams[i];
4634 int length = util_last_bit(output_usage_mask);
4635
4636 if (!(ctx->output_mask & (1ull << i)) ||
4637 output_stream != stream)
4638 continue;
4639
4640 for (unsigned j = 0; j < length; j++) {
4641 LLVMValueRef value, soffset;
4642
4643 if (!(output_usage_mask & (1 << j)))
4644 continue;
4645
4646 soffset = LLVMConstInt(ctx->ac.i32,
4647 offset *
4648 ctx->gs_max_out_vertices * 16 * 4, false);
4649
4650 offset++;
4651
4652 value = ac_build_buffer_load(&ctx->ac,
4653 ctx->gsvs_ring[0],
4654 1, ctx->ac.i32_0,
4655 vtx_offset, soffset,
4656 0, ac_glc | ac_slc, true, false);
4657
4658 LLVMTypeRef type = LLVMGetAllocatedType(ctx->abi.outputs[ac_llvm_reg_index_soa(i, j)]);
4659 if (ac_get_type_size(type) == 2) {
4660 value = LLVMBuildBitCast(ctx->ac.builder, value, ctx->ac.i32, "");
4661 value = LLVMBuildTrunc(ctx->ac.builder, value, ctx->ac.i16, "");
4662 }
4663
4664 LLVMBuildStore(ctx->ac.builder,
4665 ac_to_float(&ctx->ac, value), ctx->abi.outputs[ac_llvm_reg_index_soa(i, j)]);
4666 }
4667 }
4668
4669 if (ctx->shader_info->info.so.num_outputs)
4670 radv_emit_streamout(ctx, stream);
4671
4672 if (stream == 0) {
4673 handle_vs_outputs_post(ctx, false, true,
4674 &ctx->shader_info->vs.outinfo);
4675 }
4676
4677 LLVMBuildBr(ctx->ac.builder, end_bb);
4678 }
4679
4680 LLVMPositionBuilderAtEnd(ctx->ac.builder, end_bb);
4681 }
4682
4683 void
4684 radv_compile_gs_copy_shader(struct ac_llvm_compiler *ac_llvm,
4685 struct nir_shader *geom_shader,
4686 struct radv_shader_binary **rbinary,
4687 struct radv_shader_variant_info *shader_info,
4688 const struct radv_nir_compiler_options *options)
4689 {
4690 struct radv_shader_context ctx = {0};
4691 ctx.options = options;
4692 ctx.shader_info = shader_info;
4693
4694 enum ac_float_mode float_mode =
4695 options->unsafe_math ? AC_FLOAT_MODE_UNSAFE_FP_MATH :
4696 AC_FLOAT_MODE_DEFAULT;
4697
4698 ac_llvm_context_init(&ctx.ac, ac_llvm, options->chip_class,
4699 options->family, float_mode, 64, 64);
4700 ctx.context = ctx.ac.context;
4701
4702 ctx.is_gs_copy_shader = true;
4703 ctx.stage = MESA_SHADER_VERTEX;
4704
4705 radv_nir_shader_info_pass(geom_shader, options, &shader_info->info);
4706
4707 create_function(&ctx, MESA_SHADER_VERTEX, false, MESA_SHADER_VERTEX);
4708
4709 ctx.gs_max_out_vertices = geom_shader->info.gs.vertices_out;
4710 ac_setup_rings(&ctx);
4711
4712 nir_foreach_variable(variable, &geom_shader->outputs) {
4713 scan_shader_output_decl(&ctx, variable, geom_shader, MESA_SHADER_VERTEX);
4714 ac_handle_shader_output_decl(&ctx.ac, &ctx.abi, geom_shader,
4715 variable, MESA_SHADER_VERTEX);
4716 }
4717
4718 ac_gs_copy_shader_emit(&ctx);
4719
4720 LLVMBuildRetVoid(ctx.ac.builder);
4721
4722 ac_llvm_finalize_module(&ctx, ac_llvm->passmgr, options);
4723
4724 ac_compile_llvm_module(ac_llvm, ctx.ac.module, rbinary,
4725 MESA_SHADER_VERTEX, "GS Copy Shader", options);
4726 (*rbinary)->is_gs_copy_shader = true;
4727
4728 }