148a571fc79aa6e4e9e1a70e493bdf023232ee72
[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_info *shader_info;
52 const struct nir_shader *shader;
53 struct ac_shader_abi abi;
54
55 unsigned max_workgroup_size;
56 LLVMContextRef context;
57 LLVMValueRef main_function;
58
59 LLVMValueRef descriptor_sets[MAX_SETS];
60 LLVMValueRef ring_offsets;
61
62 LLVMValueRef vertex_buffers;
63 LLVMValueRef rel_auto_id;
64 LLVMValueRef vs_prim_id;
65 LLVMValueRef es2gs_offset;
66
67 LLVMValueRef oc_lds;
68 LLVMValueRef merged_wave_info;
69 LLVMValueRef tess_factor_offset;
70 LLVMValueRef tes_rel_patch_id;
71 LLVMValueRef tes_u;
72 LLVMValueRef tes_v;
73
74 /* HW GS */
75 /* On gfx10:
76 * - bits 0..10: ordered_wave_id
77 * - bits 12..20: number of vertices in group
78 * - bits 22..30: number of primitives in group
79 */
80 LLVMValueRef gs_tg_info;
81 LLVMValueRef gs2vs_offset;
82 LLVMValueRef gs_wave_id;
83 LLVMValueRef gs_vtx_offset[6];
84
85 LLVMValueRef esgs_ring;
86 LLVMValueRef gsvs_ring[4];
87 LLVMValueRef hs_ring_tess_offchip;
88 LLVMValueRef hs_ring_tess_factor;
89
90 /* Streamout */
91 LLVMValueRef streamout_buffers;
92 LLVMValueRef streamout_write_idx;
93 LLVMValueRef streamout_config;
94 LLVMValueRef streamout_offset[4];
95
96 gl_shader_stage stage;
97
98 LLVMValueRef inputs[RADEON_LLVM_MAX_INPUTS * 4];
99
100 uint64_t output_mask;
101
102 bool is_gs_copy_shader;
103 LLVMValueRef gs_next_vertex[4];
104 LLVMValueRef gs_curprim_verts[4];
105 LLVMValueRef gs_generated_prims[4];
106 LLVMValueRef gs_ngg_emit;
107 LLVMValueRef gs_ngg_scratch;
108
109 uint32_t tcs_num_inputs;
110 uint32_t tcs_num_patches;
111
112 LLVMValueRef vertexptr; /* GFX10 only */
113 };
114
115 struct radv_shader_output_values {
116 LLVMValueRef values[4];
117 unsigned slot_name;
118 unsigned slot_index;
119 unsigned usage_mask;
120 };
121
122 enum radeon_llvm_calling_convention {
123 RADEON_LLVM_AMDGPU_VS = 87,
124 RADEON_LLVM_AMDGPU_GS = 88,
125 RADEON_LLVM_AMDGPU_PS = 89,
126 RADEON_LLVM_AMDGPU_CS = 90,
127 RADEON_LLVM_AMDGPU_HS = 93,
128 };
129
130 static inline struct radv_shader_context *
131 radv_shader_context_from_abi(struct ac_shader_abi *abi)
132 {
133 struct radv_shader_context *ctx = NULL;
134 return container_of(abi, ctx, abi);
135 }
136
137 static LLVMValueRef get_rel_patch_id(struct radv_shader_context *ctx)
138 {
139 switch (ctx->stage) {
140 case MESA_SHADER_TESS_CTRL:
141 return ac_unpack_param(&ctx->ac, ctx->abi.tcs_rel_ids, 0, 8);
142 case MESA_SHADER_TESS_EVAL:
143 return ctx->tes_rel_patch_id;
144 break;
145 default:
146 unreachable("Illegal stage");
147 }
148 }
149
150 static unsigned
151 get_tcs_num_patches(struct radv_shader_context *ctx)
152 {
153 unsigned num_tcs_input_cp = ctx->options->key.tcs.input_vertices;
154 unsigned num_tcs_output_cp = ctx->shader->info.tess.tcs_vertices_out;
155 uint32_t input_vertex_size = ctx->tcs_num_inputs * 16;
156 uint32_t input_patch_size = ctx->options->key.tcs.input_vertices * input_vertex_size;
157 uint32_t num_tcs_outputs = util_last_bit64(ctx->shader_info->tcs.outputs_written);
158 uint32_t num_tcs_patch_outputs = util_last_bit64(ctx->shader_info->tcs.patch_outputs_written);
159 uint32_t output_vertex_size = num_tcs_outputs * 16;
160 uint32_t pervertex_output_patch_size = ctx->shader->info.tess.tcs_vertices_out * output_vertex_size;
161 uint32_t output_patch_size = pervertex_output_patch_size + num_tcs_patch_outputs * 16;
162 unsigned num_patches;
163 unsigned hardware_lds_size;
164
165 /* Ensure that we only need one wave per SIMD so we don't need to check
166 * resource usage. Also ensures that the number of tcs in and out
167 * vertices per threadgroup are at most 256.
168 */
169 num_patches = 64 / MAX2(num_tcs_input_cp, num_tcs_output_cp) * 4;
170 /* Make sure that the data fits in LDS. This assumes the shaders only
171 * use LDS for the inputs and outputs.
172 */
173 hardware_lds_size = 32768;
174
175 /* Looks like STONEY hangs if we use more than 32 KiB LDS in a single
176 * threadgroup, even though there is more than 32 KiB LDS.
177 *
178 * Test: dEQP-VK.tessellation.shader_input_output.barrier
179 */
180 if (ctx->options->chip_class >= GFX7 && ctx->options->family != CHIP_STONEY)
181 hardware_lds_size = 65536;
182
183 num_patches = MIN2(num_patches, hardware_lds_size / (input_patch_size + output_patch_size));
184 /* Make sure the output data fits in the offchip buffer */
185 num_patches = MIN2(num_patches, (ctx->options->tess_offchip_block_dw_size * 4) / output_patch_size);
186 /* Not necessary for correctness, but improves performance. The
187 * specific value is taken from the proprietary driver.
188 */
189 num_patches = MIN2(num_patches, 40);
190
191 /* GFX6 bug workaround - limit LS-HS threadgroups to only one wave. */
192 if (ctx->options->chip_class == GFX6) {
193 unsigned one_wave = 64 / MAX2(num_tcs_input_cp, num_tcs_output_cp);
194 num_patches = MIN2(num_patches, one_wave);
195 }
196 return num_patches;
197 }
198
199 static unsigned
200 calculate_tess_lds_size(struct radv_shader_context *ctx)
201 {
202 unsigned num_tcs_input_cp = ctx->options->key.tcs.input_vertices;
203 unsigned num_tcs_output_cp;
204 unsigned num_tcs_outputs, num_tcs_patch_outputs;
205 unsigned input_vertex_size, output_vertex_size;
206 unsigned input_patch_size, output_patch_size;
207 unsigned pervertex_output_patch_size;
208 unsigned output_patch0_offset;
209 unsigned num_patches;
210 unsigned lds_size;
211
212 num_tcs_output_cp = ctx->shader->info.tess.tcs_vertices_out;
213 num_tcs_outputs = util_last_bit64(ctx->shader_info->tcs.outputs_written);
214 num_tcs_patch_outputs = util_last_bit64(ctx->shader_info->tcs.patch_outputs_written);
215
216 input_vertex_size = ctx->tcs_num_inputs * 16;
217 output_vertex_size = num_tcs_outputs * 16;
218
219 input_patch_size = num_tcs_input_cp * input_vertex_size;
220
221 pervertex_output_patch_size = num_tcs_output_cp * output_vertex_size;
222 output_patch_size = pervertex_output_patch_size + num_tcs_patch_outputs * 16;
223
224 num_patches = ctx->tcs_num_patches;
225 output_patch0_offset = input_patch_size * num_patches;
226
227 lds_size = output_patch0_offset + output_patch_size * num_patches;
228 return lds_size;
229 }
230
231 /* Tessellation shaders pass outputs to the next shader using LDS.
232 *
233 * LS outputs = TCS inputs
234 * TCS outputs = TES inputs
235 *
236 * The LDS layout is:
237 * - TCS inputs for patch 0
238 * - TCS inputs for patch 1
239 * - TCS inputs for patch 2 = get_tcs_in_current_patch_offset (if RelPatchID==2)
240 * - ...
241 * - TCS outputs for patch 0 = get_tcs_out_patch0_offset
242 * - Per-patch TCS outputs for patch 0 = get_tcs_out_patch0_patch_data_offset
243 * - TCS outputs for patch 1
244 * - Per-patch TCS outputs for patch 1
245 * - TCS outputs for patch 2 = get_tcs_out_current_patch_offset (if RelPatchID==2)
246 * - Per-patch TCS outputs for patch 2 = get_tcs_out_current_patch_data_offset (if RelPatchID==2)
247 * - ...
248 *
249 * All three shaders VS(LS), TCS, TES share the same LDS space.
250 */
251 static LLVMValueRef
252 get_tcs_in_patch_stride(struct radv_shader_context *ctx)
253 {
254 assert (ctx->stage == MESA_SHADER_TESS_CTRL);
255 uint32_t input_vertex_size = ctx->tcs_num_inputs * 16;
256 uint32_t input_patch_size = ctx->options->key.tcs.input_vertices * input_vertex_size;
257
258 input_patch_size /= 4;
259 return LLVMConstInt(ctx->ac.i32, input_patch_size, false);
260 }
261
262 static LLVMValueRef
263 get_tcs_out_patch_stride(struct radv_shader_context *ctx)
264 {
265 uint32_t num_tcs_outputs = util_last_bit64(ctx->shader_info->tcs.outputs_written);
266 uint32_t num_tcs_patch_outputs = util_last_bit64(ctx->shader_info->tcs.patch_outputs_written);
267 uint32_t output_vertex_size = num_tcs_outputs * 16;
268 uint32_t pervertex_output_patch_size = ctx->shader->info.tess.tcs_vertices_out * output_vertex_size;
269 uint32_t output_patch_size = pervertex_output_patch_size + num_tcs_patch_outputs * 16;
270 output_patch_size /= 4;
271 return LLVMConstInt(ctx->ac.i32, output_patch_size, false);
272 }
273
274 static LLVMValueRef
275 get_tcs_out_vertex_stride(struct radv_shader_context *ctx)
276 {
277 uint32_t num_tcs_outputs = util_last_bit64(ctx->shader_info->tcs.outputs_written);
278 uint32_t output_vertex_size = num_tcs_outputs * 16;
279 output_vertex_size /= 4;
280 return LLVMConstInt(ctx->ac.i32, output_vertex_size, false);
281 }
282
283 static LLVMValueRef
284 get_tcs_out_patch0_offset(struct radv_shader_context *ctx)
285 {
286 assert (ctx->stage == MESA_SHADER_TESS_CTRL);
287 uint32_t input_vertex_size = ctx->tcs_num_inputs * 16;
288 uint32_t input_patch_size = ctx->options->key.tcs.input_vertices * input_vertex_size;
289 uint32_t output_patch0_offset = input_patch_size;
290 unsigned num_patches = ctx->tcs_num_patches;
291
292 output_patch0_offset *= num_patches;
293 output_patch0_offset /= 4;
294 return LLVMConstInt(ctx->ac.i32, output_patch0_offset, false);
295 }
296
297 static LLVMValueRef
298 get_tcs_out_patch0_patch_data_offset(struct radv_shader_context *ctx)
299 {
300 assert (ctx->stage == MESA_SHADER_TESS_CTRL);
301 uint32_t input_vertex_size = ctx->tcs_num_inputs * 16;
302 uint32_t input_patch_size = ctx->options->key.tcs.input_vertices * input_vertex_size;
303 uint32_t output_patch0_offset = input_patch_size;
304
305 uint32_t num_tcs_outputs = util_last_bit64(ctx->shader_info->tcs.outputs_written);
306 uint32_t output_vertex_size = num_tcs_outputs * 16;
307 uint32_t pervertex_output_patch_size = ctx->shader->info.tess.tcs_vertices_out * output_vertex_size;
308 unsigned num_patches = ctx->tcs_num_patches;
309
310 output_patch0_offset *= num_patches;
311 output_patch0_offset += pervertex_output_patch_size;
312 output_patch0_offset /= 4;
313 return LLVMConstInt(ctx->ac.i32, output_patch0_offset, false);
314 }
315
316 static LLVMValueRef
317 get_tcs_in_current_patch_offset(struct radv_shader_context *ctx)
318 {
319 LLVMValueRef patch_stride = get_tcs_in_patch_stride(ctx);
320 LLVMValueRef rel_patch_id = get_rel_patch_id(ctx);
321
322 return LLVMBuildMul(ctx->ac.builder, patch_stride, rel_patch_id, "");
323 }
324
325 static LLVMValueRef
326 get_tcs_out_current_patch_offset(struct radv_shader_context *ctx)
327 {
328 LLVMValueRef patch0_offset = get_tcs_out_patch0_offset(ctx);
329 LLVMValueRef patch_stride = get_tcs_out_patch_stride(ctx);
330 LLVMValueRef rel_patch_id = get_rel_patch_id(ctx);
331
332 return ac_build_imad(&ctx->ac, patch_stride, rel_patch_id,
333 patch0_offset);
334 }
335
336 static LLVMValueRef
337 get_tcs_out_current_patch_data_offset(struct radv_shader_context *ctx)
338 {
339 LLVMValueRef patch0_patch_data_offset =
340 get_tcs_out_patch0_patch_data_offset(ctx);
341 LLVMValueRef patch_stride = get_tcs_out_patch_stride(ctx);
342 LLVMValueRef rel_patch_id = get_rel_patch_id(ctx);
343
344 return ac_build_imad(&ctx->ac, patch_stride, rel_patch_id,
345 patch0_patch_data_offset);
346 }
347
348 #define MAX_ARGS 64
349 struct arg_info {
350 LLVMTypeRef types[MAX_ARGS];
351 LLVMValueRef *assign[MAX_ARGS];
352 uint8_t count;
353 uint8_t sgpr_count;
354 uint8_t num_sgprs_used;
355 uint8_t num_vgprs_used;
356 };
357
358 enum radv_arg_regfile {
359 ARG_SGPR,
360 ARG_VGPR,
361 };
362
363 static void
364 add_arg(struct arg_info *info, enum radv_arg_regfile regfile, LLVMTypeRef type,
365 LLVMValueRef *param_ptr)
366 {
367 assert(info->count < MAX_ARGS);
368
369 info->assign[info->count] = param_ptr;
370 info->types[info->count] = type;
371 info->count++;
372
373 if (regfile == ARG_SGPR) {
374 info->num_sgprs_used += ac_get_type_size(type) / 4;
375 info->sgpr_count++;
376 } else {
377 assert(regfile == ARG_VGPR);
378 info->num_vgprs_used += ac_get_type_size(type) / 4;
379 }
380 }
381
382 static void assign_arguments(LLVMValueRef main_function,
383 struct arg_info *info)
384 {
385 unsigned i;
386 for (i = 0; i < info->count; i++) {
387 if (info->assign[i])
388 *info->assign[i] = LLVMGetParam(main_function, i);
389 }
390 }
391
392 static LLVMValueRef
393 create_llvm_function(LLVMContextRef ctx, LLVMModuleRef module,
394 LLVMBuilderRef builder, LLVMTypeRef *return_types,
395 unsigned num_return_elems,
396 struct arg_info *args,
397 unsigned max_workgroup_size,
398 const struct radv_nir_compiler_options *options)
399 {
400 LLVMTypeRef main_function_type, ret_type;
401 LLVMBasicBlockRef main_function_body;
402
403 if (num_return_elems)
404 ret_type = LLVMStructTypeInContext(ctx, return_types,
405 num_return_elems, true);
406 else
407 ret_type = LLVMVoidTypeInContext(ctx);
408
409 /* Setup the function */
410 main_function_type =
411 LLVMFunctionType(ret_type, args->types, args->count, 0);
412 LLVMValueRef main_function =
413 LLVMAddFunction(module, "main", main_function_type);
414 main_function_body =
415 LLVMAppendBasicBlockInContext(ctx, main_function, "main_body");
416 LLVMPositionBuilderAtEnd(builder, main_function_body);
417
418 LLVMSetFunctionCallConv(main_function, RADEON_LLVM_AMDGPU_CS);
419 for (unsigned i = 0; i < args->sgpr_count; ++i) {
420 LLVMValueRef P = LLVMGetParam(main_function, i);
421
422 ac_add_function_attr(ctx, main_function, i + 1, AC_FUNC_ATTR_INREG);
423
424 if (LLVMGetTypeKind(LLVMTypeOf(P)) == LLVMPointerTypeKind) {
425 ac_add_function_attr(ctx, main_function, i + 1, AC_FUNC_ATTR_NOALIAS);
426 ac_add_attr_dereferenceable(P, UINT64_MAX);
427 }
428 }
429
430 if (options->address32_hi) {
431 ac_llvm_add_target_dep_function_attr(main_function,
432 "amdgpu-32bit-address-high-bits",
433 options->address32_hi);
434 }
435
436 ac_llvm_set_workgroup_size(main_function, max_workgroup_size);
437
438 return main_function;
439 }
440
441
442 static void
443 set_loc(struct radv_userdata_info *ud_info, uint8_t *sgpr_idx,
444 uint8_t num_sgprs)
445 {
446 ud_info->sgpr_idx = *sgpr_idx;
447 ud_info->num_sgprs = num_sgprs;
448 *sgpr_idx += num_sgprs;
449 }
450
451 static void
452 set_loc_shader(struct radv_shader_context *ctx, int idx, uint8_t *sgpr_idx,
453 uint8_t num_sgprs)
454 {
455 struct radv_userdata_info *ud_info =
456 &ctx->shader_info->user_sgprs_locs.shader_data[idx];
457 assert(ud_info);
458
459 set_loc(ud_info, sgpr_idx, num_sgprs);
460 }
461
462 static void
463 set_loc_shader_ptr(struct radv_shader_context *ctx, int idx, uint8_t *sgpr_idx)
464 {
465 bool use_32bit_pointers = idx != AC_UD_SCRATCH_RING_OFFSETS;
466
467 set_loc_shader(ctx, idx, sgpr_idx, use_32bit_pointers ? 1 : 2);
468 }
469
470 static void
471 set_loc_desc(struct radv_shader_context *ctx, int idx, uint8_t *sgpr_idx)
472 {
473 struct radv_userdata_locations *locs =
474 &ctx->shader_info->user_sgprs_locs;
475 struct radv_userdata_info *ud_info = &locs->descriptor_sets[idx];
476 assert(ud_info);
477
478 set_loc(ud_info, sgpr_idx, 1);
479
480 locs->descriptor_sets_enabled |= 1 << idx;
481 }
482
483 struct user_sgpr_info {
484 bool need_ring_offsets;
485 bool indirect_all_descriptor_sets;
486 uint8_t remaining_sgprs;
487 };
488
489 static bool needs_view_index_sgpr(struct radv_shader_context *ctx,
490 gl_shader_stage stage)
491 {
492 switch (stage) {
493 case MESA_SHADER_VERTEX:
494 if (ctx->shader_info->needs_multiview_view_index ||
495 (!ctx->options->key.vs_common_out.as_es && !ctx->options->key.vs_common_out.as_ls && ctx->options->key.has_multiview_view_index))
496 return true;
497 break;
498 case MESA_SHADER_TESS_EVAL:
499 if (ctx->shader_info->needs_multiview_view_index || (!ctx->options->key.vs_common_out.as_es && ctx->options->key.has_multiview_view_index))
500 return true;
501 break;
502 case MESA_SHADER_GEOMETRY:
503 case MESA_SHADER_TESS_CTRL:
504 if (ctx->shader_info->needs_multiview_view_index)
505 return true;
506 break;
507 default:
508 break;
509 }
510 return false;
511 }
512
513 static uint8_t
514 count_vs_user_sgprs(struct radv_shader_context *ctx)
515 {
516 uint8_t count = 0;
517
518 if (ctx->shader_info->vs.has_vertex_buffers)
519 count++;
520 count += ctx->shader_info->vs.needs_draw_id ? 3 : 2;
521
522 return count;
523 }
524
525 static void allocate_inline_push_consts(struct radv_shader_context *ctx,
526 struct user_sgpr_info *user_sgpr_info)
527 {
528 uint8_t remaining_sgprs = user_sgpr_info->remaining_sgprs;
529
530 /* Only supported if shaders use push constants. */
531 if (ctx->shader_info->min_push_constant_used == UINT8_MAX)
532 return;
533
534 /* Only supported if shaders don't have indirect push constants. */
535 if (ctx->shader_info->has_indirect_push_constants)
536 return;
537
538 /* Only supported for 32-bit push constants. */
539 if (!ctx->shader_info->has_only_32bit_push_constants)
540 return;
541
542 uint8_t num_push_consts =
543 (ctx->shader_info->max_push_constant_used -
544 ctx->shader_info->min_push_constant_used) / 4;
545
546 /* Check if the number of user SGPRs is large enough. */
547 if (num_push_consts < remaining_sgprs) {
548 ctx->shader_info->num_inline_push_consts = num_push_consts;
549 } else {
550 ctx->shader_info->num_inline_push_consts = remaining_sgprs;
551 }
552
553 /* Clamp to the maximum number of allowed inlined push constants. */
554 if (ctx->shader_info->num_inline_push_consts > AC_MAX_INLINE_PUSH_CONSTS)
555 ctx->shader_info->num_inline_push_consts = AC_MAX_INLINE_PUSH_CONSTS;
556
557 if (ctx->shader_info->num_inline_push_consts == num_push_consts &&
558 !ctx->shader_info->loads_dynamic_offsets) {
559 /* Disable the default push constants path if all constants are
560 * inlined and if shaders don't use dynamic descriptors.
561 */
562 ctx->shader_info->loads_push_constants = false;
563 }
564
565 ctx->shader_info->base_inline_push_consts =
566 ctx->shader_info->min_push_constant_used / 4;
567 }
568
569 static void allocate_user_sgprs(struct radv_shader_context *ctx,
570 gl_shader_stage stage,
571 bool has_previous_stage,
572 gl_shader_stage previous_stage,
573 bool needs_view_index,
574 struct user_sgpr_info *user_sgpr_info)
575 {
576 uint8_t user_sgpr_count = 0;
577
578 memset(user_sgpr_info, 0, sizeof(struct user_sgpr_info));
579
580 /* until we sort out scratch/global buffers always assign ring offsets for gs/vs/es */
581 if (stage == MESA_SHADER_GEOMETRY ||
582 stage == MESA_SHADER_VERTEX ||
583 stage == MESA_SHADER_TESS_CTRL ||
584 stage == MESA_SHADER_TESS_EVAL ||
585 ctx->is_gs_copy_shader)
586 user_sgpr_info->need_ring_offsets = true;
587
588 if (stage == MESA_SHADER_FRAGMENT &&
589 ctx->shader_info->ps.needs_sample_positions)
590 user_sgpr_info->need_ring_offsets = true;
591
592 /* 2 user sgprs will nearly always be allocated for scratch/rings */
593 if (ctx->options->supports_spill || user_sgpr_info->need_ring_offsets) {
594 user_sgpr_count += 2;
595 }
596
597 switch (stage) {
598 case MESA_SHADER_COMPUTE:
599 if (ctx->shader_info->cs.uses_grid_size)
600 user_sgpr_count += 3;
601 break;
602 case MESA_SHADER_FRAGMENT:
603 user_sgpr_count += ctx->shader_info->ps.needs_sample_positions;
604 break;
605 case MESA_SHADER_VERTEX:
606 if (!ctx->is_gs_copy_shader)
607 user_sgpr_count += count_vs_user_sgprs(ctx);
608 break;
609 case MESA_SHADER_TESS_CTRL:
610 if (has_previous_stage) {
611 if (previous_stage == MESA_SHADER_VERTEX)
612 user_sgpr_count += count_vs_user_sgprs(ctx);
613 }
614 break;
615 case MESA_SHADER_TESS_EVAL:
616 break;
617 case MESA_SHADER_GEOMETRY:
618 if (has_previous_stage) {
619 if (previous_stage == MESA_SHADER_VERTEX) {
620 user_sgpr_count += count_vs_user_sgprs(ctx);
621 }
622 }
623 break;
624 default:
625 break;
626 }
627
628 if (needs_view_index)
629 user_sgpr_count++;
630
631 if (ctx->shader_info->loads_push_constants)
632 user_sgpr_count++;
633
634 if (ctx->shader_info->so.num_outputs)
635 user_sgpr_count++;
636
637 uint32_t available_sgprs = ctx->options->chip_class >= GFX9 && stage != MESA_SHADER_COMPUTE ? 32 : 16;
638 uint32_t remaining_sgprs = available_sgprs - user_sgpr_count;
639 uint32_t num_desc_set =
640 util_bitcount(ctx->shader_info->desc_set_used_mask);
641
642 if (remaining_sgprs < num_desc_set) {
643 user_sgpr_info->indirect_all_descriptor_sets = true;
644 user_sgpr_info->remaining_sgprs = remaining_sgprs - 1;
645 } else {
646 user_sgpr_info->remaining_sgprs = remaining_sgprs - num_desc_set;
647 }
648
649 allocate_inline_push_consts(ctx, user_sgpr_info);
650 }
651
652 static void
653 declare_global_input_sgprs(struct radv_shader_context *ctx,
654 const struct user_sgpr_info *user_sgpr_info,
655 struct arg_info *args,
656 LLVMValueRef *desc_sets)
657 {
658 LLVMTypeRef type = ac_array_in_const32_addr_space(ctx->ac.i8);
659
660 /* 1 for each descriptor set */
661 if (!user_sgpr_info->indirect_all_descriptor_sets) {
662 uint32_t mask = ctx->shader_info->desc_set_used_mask;
663
664 while (mask) {
665 int i = u_bit_scan(&mask);
666
667 add_arg(args, ARG_SGPR, type, &ctx->descriptor_sets[i]);
668 }
669 } else {
670 add_arg(args, ARG_SGPR, ac_array_in_const32_addr_space(type),
671 desc_sets);
672 }
673
674 if (ctx->shader_info->loads_push_constants) {
675 /* 1 for push constants and dynamic descriptors */
676 add_arg(args, ARG_SGPR, type, &ctx->abi.push_constants);
677 }
678
679 for (unsigned i = 0; i < ctx->shader_info->num_inline_push_consts; i++) {
680 add_arg(args, ARG_SGPR, ctx->ac.i32,
681 &ctx->abi.inline_push_consts[i]);
682 }
683 ctx->abi.num_inline_push_consts = ctx->shader_info->num_inline_push_consts;
684 ctx->abi.base_inline_push_consts = ctx->shader_info->base_inline_push_consts;
685
686 if (ctx->shader_info->so.num_outputs) {
687 add_arg(args, ARG_SGPR,
688 ac_array_in_const32_addr_space(ctx->ac.v4i32),
689 &ctx->streamout_buffers);
690 }
691 }
692
693 static void
694 declare_vs_specific_input_sgprs(struct radv_shader_context *ctx,
695 gl_shader_stage stage,
696 bool has_previous_stage,
697 gl_shader_stage previous_stage,
698 struct arg_info *args)
699 {
700 if (!ctx->is_gs_copy_shader &&
701 (stage == MESA_SHADER_VERTEX ||
702 (has_previous_stage && previous_stage == MESA_SHADER_VERTEX))) {
703 if (ctx->shader_info->vs.has_vertex_buffers) {
704 add_arg(args, ARG_SGPR,
705 ac_array_in_const32_addr_space(ctx->ac.v4i32),
706 &ctx->vertex_buffers);
707 }
708 add_arg(args, ARG_SGPR, ctx->ac.i32, &ctx->abi.base_vertex);
709 add_arg(args, ARG_SGPR, ctx->ac.i32, &ctx->abi.start_instance);
710 if (ctx->shader_info->vs.needs_draw_id) {
711 add_arg(args, ARG_SGPR, ctx->ac.i32, &ctx->abi.draw_id);
712 }
713 }
714 }
715
716 static void
717 declare_vs_input_vgprs(struct radv_shader_context *ctx, struct arg_info *args)
718 {
719 add_arg(args, ARG_VGPR, ctx->ac.i32, &ctx->abi.vertex_id);
720 if (!ctx->is_gs_copy_shader) {
721 if (ctx->options->key.vs_common_out.as_ls) {
722 add_arg(args, ARG_VGPR, ctx->ac.i32, &ctx->rel_auto_id);
723 if (ctx->ac.chip_class >= GFX10) {
724 add_arg(args, ARG_VGPR, ctx->ac.i32, NULL); /* user vgpr */
725 add_arg(args, ARG_VGPR, ctx->ac.i32, &ctx->abi.instance_id);
726 } else {
727 add_arg(args, ARG_VGPR, ctx->ac.i32, &ctx->abi.instance_id);
728 add_arg(args, ARG_VGPR, ctx->ac.i32, NULL); /* unused */
729 }
730 } else {
731 if (ctx->ac.chip_class >= GFX10) {
732 if (ctx->options->key.vs_common_out.as_ngg) {
733 add_arg(args, ARG_VGPR, ctx->ac.i32, NULL); /* user vgpr */
734 add_arg(args, ARG_VGPR, ctx->ac.i32, NULL); /* user vgpr */
735 add_arg(args, ARG_VGPR, ctx->ac.i32, &ctx->abi.instance_id);
736 } else {
737 add_arg(args, ARG_VGPR, ctx->ac.i32, NULL); /* unused */
738 add_arg(args, ARG_VGPR, ctx->ac.i32, &ctx->vs_prim_id);
739 add_arg(args, ARG_VGPR, ctx->ac.i32, &ctx->abi.instance_id);
740 }
741 } else {
742 add_arg(args, ARG_VGPR, ctx->ac.i32, &ctx->abi.instance_id);
743 add_arg(args, ARG_VGPR, ctx->ac.i32, &ctx->vs_prim_id);
744 add_arg(args, ARG_VGPR, ctx->ac.i32, NULL); /* unused */
745 }
746 }
747 }
748 }
749
750 static void
751 declare_streamout_sgprs(struct radv_shader_context *ctx, gl_shader_stage stage,
752 struct arg_info *args)
753 {
754 int i;
755
756 if (ctx->options->use_ngg_streamout)
757 return;
758
759 /* Streamout SGPRs. */
760 if (ctx->shader_info->so.num_outputs) {
761 assert(stage == MESA_SHADER_VERTEX ||
762 stage == MESA_SHADER_TESS_EVAL);
763
764 if (stage != MESA_SHADER_TESS_EVAL) {
765 add_arg(args, ARG_SGPR, ctx->ac.i32, &ctx->streamout_config);
766 } else {
767 args->assign[args->count - 1] = &ctx->streamout_config;
768 args->types[args->count - 1] = ctx->ac.i32;
769 }
770
771 add_arg(args, ARG_SGPR, ctx->ac.i32, &ctx->streamout_write_idx);
772 }
773
774 /* A streamout buffer offset is loaded if the stride is non-zero. */
775 for (i = 0; i < 4; i++) {
776 if (!ctx->shader_info->so.strides[i])
777 continue;
778
779 add_arg(args, ARG_SGPR, ctx->ac.i32, &ctx->streamout_offset[i]);
780 }
781 }
782
783 static void
784 declare_tes_input_vgprs(struct radv_shader_context *ctx, struct arg_info *args)
785 {
786 add_arg(args, ARG_VGPR, ctx->ac.f32, &ctx->tes_u);
787 add_arg(args, ARG_VGPR, ctx->ac.f32, &ctx->tes_v);
788 add_arg(args, ARG_VGPR, ctx->ac.i32, &ctx->tes_rel_patch_id);
789 add_arg(args, ARG_VGPR, ctx->ac.i32, &ctx->abi.tes_patch_id);
790 }
791
792 static void
793 set_global_input_locs(struct radv_shader_context *ctx,
794 const struct user_sgpr_info *user_sgpr_info,
795 LLVMValueRef desc_sets, uint8_t *user_sgpr_idx)
796 {
797 uint32_t mask = ctx->shader_info->desc_set_used_mask;
798
799 if (!user_sgpr_info->indirect_all_descriptor_sets) {
800 while (mask) {
801 int i = u_bit_scan(&mask);
802
803 set_loc_desc(ctx, i, user_sgpr_idx);
804 }
805 } else {
806 set_loc_shader_ptr(ctx, AC_UD_INDIRECT_DESCRIPTOR_SETS,
807 user_sgpr_idx);
808
809 while (mask) {
810 int i = u_bit_scan(&mask);
811
812 ctx->descriptor_sets[i] =
813 ac_build_load_to_sgpr(&ctx->ac, desc_sets,
814 LLVMConstInt(ctx->ac.i32, i, false));
815
816 }
817
818 ctx->shader_info->need_indirect_descriptor_sets = true;
819 }
820
821 if (ctx->shader_info->loads_push_constants) {
822 set_loc_shader_ptr(ctx, AC_UD_PUSH_CONSTANTS, user_sgpr_idx);
823 }
824
825 if (ctx->shader_info->num_inline_push_consts) {
826 set_loc_shader(ctx, AC_UD_INLINE_PUSH_CONSTANTS, user_sgpr_idx,
827 ctx->shader_info->num_inline_push_consts);
828 }
829
830 if (ctx->streamout_buffers) {
831 set_loc_shader_ptr(ctx, AC_UD_STREAMOUT_BUFFERS,
832 user_sgpr_idx);
833 }
834 }
835
836 static void
837 set_vs_specific_input_locs(struct radv_shader_context *ctx,
838 gl_shader_stage stage, bool has_previous_stage,
839 gl_shader_stage previous_stage,
840 uint8_t *user_sgpr_idx)
841 {
842 if (!ctx->is_gs_copy_shader &&
843 (stage == MESA_SHADER_VERTEX ||
844 (has_previous_stage && previous_stage == MESA_SHADER_VERTEX))) {
845 if (ctx->shader_info->vs.has_vertex_buffers) {
846 set_loc_shader_ptr(ctx, AC_UD_VS_VERTEX_BUFFERS,
847 user_sgpr_idx);
848 }
849
850 unsigned vs_num = 2;
851 if (ctx->shader_info->vs.needs_draw_id)
852 vs_num++;
853
854 set_loc_shader(ctx, AC_UD_VS_BASE_VERTEX_START_INSTANCE,
855 user_sgpr_idx, vs_num);
856 }
857 }
858
859 static void set_llvm_calling_convention(LLVMValueRef func,
860 gl_shader_stage stage)
861 {
862 enum radeon_llvm_calling_convention calling_conv;
863
864 switch (stage) {
865 case MESA_SHADER_VERTEX:
866 case MESA_SHADER_TESS_EVAL:
867 calling_conv = RADEON_LLVM_AMDGPU_VS;
868 break;
869 case MESA_SHADER_GEOMETRY:
870 calling_conv = RADEON_LLVM_AMDGPU_GS;
871 break;
872 case MESA_SHADER_TESS_CTRL:
873 calling_conv = RADEON_LLVM_AMDGPU_HS;
874 break;
875 case MESA_SHADER_FRAGMENT:
876 calling_conv = RADEON_LLVM_AMDGPU_PS;
877 break;
878 case MESA_SHADER_COMPUTE:
879 calling_conv = RADEON_LLVM_AMDGPU_CS;
880 break;
881 default:
882 unreachable("Unhandle shader type");
883 }
884
885 LLVMSetFunctionCallConv(func, calling_conv);
886 }
887
888 /* Returns whether the stage is a stage that can be directly before the GS */
889 static bool is_pre_gs_stage(gl_shader_stage stage)
890 {
891 return stage == MESA_SHADER_VERTEX || stage == MESA_SHADER_TESS_EVAL;
892 }
893
894 static void create_function(struct radv_shader_context *ctx,
895 gl_shader_stage stage,
896 bool has_previous_stage,
897 gl_shader_stage previous_stage)
898 {
899 uint8_t user_sgpr_idx;
900 struct user_sgpr_info user_sgpr_info;
901 struct arg_info args = {};
902 LLVMValueRef desc_sets;
903 bool needs_view_index = needs_view_index_sgpr(ctx, stage);
904
905 if (ctx->ac.chip_class >= GFX10) {
906 if (is_pre_gs_stage(stage) && ctx->options->key.vs_common_out.as_ngg) {
907 /* On GFX10, VS is merged into GS for NGG. */
908 previous_stage = stage;
909 stage = MESA_SHADER_GEOMETRY;
910 has_previous_stage = true;
911 }
912 }
913
914 allocate_user_sgprs(ctx, stage, has_previous_stage,
915 previous_stage, needs_view_index, &user_sgpr_info);
916
917 if (user_sgpr_info.need_ring_offsets && !ctx->options->supports_spill) {
918 add_arg(&args, ARG_SGPR, ac_array_in_const_addr_space(ctx->ac.v4i32),
919 &ctx->ring_offsets);
920 }
921
922 switch (stage) {
923 case MESA_SHADER_COMPUTE:
924 declare_global_input_sgprs(ctx, &user_sgpr_info, &args,
925 &desc_sets);
926
927 if (ctx->shader_info->cs.uses_grid_size) {
928 add_arg(&args, ARG_SGPR, ctx->ac.v3i32,
929 &ctx->abi.num_work_groups);
930 }
931
932 for (int i = 0; i < 3; i++) {
933 ctx->abi.workgroup_ids[i] = NULL;
934 if (ctx->shader_info->cs.uses_block_id[i]) {
935 add_arg(&args, ARG_SGPR, ctx->ac.i32,
936 &ctx->abi.workgroup_ids[i]);
937 }
938 }
939
940 if (ctx->shader_info->cs.uses_local_invocation_idx)
941 add_arg(&args, ARG_SGPR, ctx->ac.i32, &ctx->abi.tg_size);
942 add_arg(&args, ARG_VGPR, ctx->ac.v3i32,
943 &ctx->abi.local_invocation_ids);
944 break;
945 case MESA_SHADER_VERTEX:
946 declare_global_input_sgprs(ctx, &user_sgpr_info, &args,
947 &desc_sets);
948
949 declare_vs_specific_input_sgprs(ctx, stage, has_previous_stage,
950 previous_stage, &args);
951
952 if (needs_view_index)
953 add_arg(&args, ARG_SGPR, ctx->ac.i32,
954 &ctx->abi.view_index);
955 if (ctx->options->key.vs_common_out.as_es) {
956 add_arg(&args, ARG_SGPR, ctx->ac.i32,
957 &ctx->es2gs_offset);
958 } else if (ctx->options->key.vs_common_out.as_ls) {
959 /* no extra parameters */
960 } else {
961 declare_streamout_sgprs(ctx, stage, &args);
962 }
963
964 declare_vs_input_vgprs(ctx, &args);
965 break;
966 case MESA_SHADER_TESS_CTRL:
967 if (has_previous_stage) {
968 // First 6 system regs
969 add_arg(&args, ARG_SGPR, ctx->ac.i32, &ctx->oc_lds);
970 add_arg(&args, ARG_SGPR, ctx->ac.i32,
971 &ctx->merged_wave_info);
972 add_arg(&args, ARG_SGPR, ctx->ac.i32,
973 &ctx->tess_factor_offset);
974
975 add_arg(&args, ARG_SGPR, ctx->ac.i32, NULL); // scratch offset
976 add_arg(&args, ARG_SGPR, ctx->ac.i32, NULL); // unknown
977 add_arg(&args, ARG_SGPR, ctx->ac.i32, NULL); // unknown
978
979 declare_global_input_sgprs(ctx, &user_sgpr_info, &args,
980 &desc_sets);
981
982 declare_vs_specific_input_sgprs(ctx, stage,
983 has_previous_stage,
984 previous_stage, &args);
985
986 if (needs_view_index)
987 add_arg(&args, ARG_SGPR, ctx->ac.i32,
988 &ctx->abi.view_index);
989
990 add_arg(&args, ARG_VGPR, ctx->ac.i32,
991 &ctx->abi.tcs_patch_id);
992 add_arg(&args, ARG_VGPR, ctx->ac.i32,
993 &ctx->abi.tcs_rel_ids);
994
995 declare_vs_input_vgprs(ctx, &args);
996 } else {
997 declare_global_input_sgprs(ctx, &user_sgpr_info, &args,
998 &desc_sets);
999
1000 if (needs_view_index)
1001 add_arg(&args, ARG_SGPR, ctx->ac.i32,
1002 &ctx->abi.view_index);
1003
1004 add_arg(&args, ARG_SGPR, ctx->ac.i32, &ctx->oc_lds);
1005 add_arg(&args, ARG_SGPR, ctx->ac.i32,
1006 &ctx->tess_factor_offset);
1007 add_arg(&args, ARG_VGPR, ctx->ac.i32,
1008 &ctx->abi.tcs_patch_id);
1009 add_arg(&args, ARG_VGPR, ctx->ac.i32,
1010 &ctx->abi.tcs_rel_ids);
1011 }
1012 break;
1013 case MESA_SHADER_TESS_EVAL:
1014 declare_global_input_sgprs(ctx, &user_sgpr_info, &args,
1015 &desc_sets);
1016
1017 if (needs_view_index)
1018 add_arg(&args, ARG_SGPR, ctx->ac.i32,
1019 &ctx->abi.view_index);
1020
1021 if (ctx->options->key.vs_common_out.as_es) {
1022 add_arg(&args, ARG_SGPR, ctx->ac.i32, &ctx->oc_lds);
1023 add_arg(&args, ARG_SGPR, ctx->ac.i32, NULL);
1024 add_arg(&args, ARG_SGPR, ctx->ac.i32,
1025 &ctx->es2gs_offset);
1026 } else {
1027 add_arg(&args, ARG_SGPR, ctx->ac.i32, NULL);
1028 declare_streamout_sgprs(ctx, stage, &args);
1029 add_arg(&args, ARG_SGPR, ctx->ac.i32, &ctx->oc_lds);
1030 }
1031 declare_tes_input_vgprs(ctx, &args);
1032 break;
1033 case MESA_SHADER_GEOMETRY:
1034 if (has_previous_stage) {
1035 // First 6 system regs
1036 if (ctx->options->key.vs_common_out.as_ngg) {
1037 add_arg(&args, ARG_SGPR, ctx->ac.i32,
1038 &ctx->gs_tg_info);
1039 } else {
1040 add_arg(&args, ARG_SGPR, ctx->ac.i32,
1041 &ctx->gs2vs_offset);
1042 }
1043
1044 add_arg(&args, ARG_SGPR, ctx->ac.i32,
1045 &ctx->merged_wave_info);
1046 add_arg(&args, ARG_SGPR, ctx->ac.i32, &ctx->oc_lds);
1047
1048 add_arg(&args, ARG_SGPR, ctx->ac.i32, NULL); // scratch offset
1049 add_arg(&args, ARG_SGPR, ctx->ac.i32, NULL); // unknown
1050 add_arg(&args, ARG_SGPR, ctx->ac.i32, NULL); // unknown
1051
1052 declare_global_input_sgprs(ctx, &user_sgpr_info, &args,
1053 &desc_sets);
1054
1055 if (previous_stage != MESA_SHADER_TESS_EVAL) {
1056 declare_vs_specific_input_sgprs(ctx, stage,
1057 has_previous_stage,
1058 previous_stage,
1059 &args);
1060 }
1061
1062 if (needs_view_index)
1063 add_arg(&args, ARG_SGPR, ctx->ac.i32,
1064 &ctx->abi.view_index);
1065
1066 add_arg(&args, ARG_VGPR, ctx->ac.i32,
1067 &ctx->gs_vtx_offset[0]);
1068 add_arg(&args, ARG_VGPR, ctx->ac.i32,
1069 &ctx->gs_vtx_offset[2]);
1070 add_arg(&args, ARG_VGPR, ctx->ac.i32,
1071 &ctx->abi.gs_prim_id);
1072 add_arg(&args, ARG_VGPR, ctx->ac.i32,
1073 &ctx->abi.gs_invocation_id);
1074 add_arg(&args, ARG_VGPR, ctx->ac.i32,
1075 &ctx->gs_vtx_offset[4]);
1076
1077 if (previous_stage == MESA_SHADER_VERTEX) {
1078 declare_vs_input_vgprs(ctx, &args);
1079 } else {
1080 declare_tes_input_vgprs(ctx, &args);
1081 }
1082 } else {
1083 declare_global_input_sgprs(ctx, &user_sgpr_info, &args,
1084 &desc_sets);
1085
1086 if (needs_view_index)
1087 add_arg(&args, ARG_SGPR, ctx->ac.i32,
1088 &ctx->abi.view_index);
1089
1090 add_arg(&args, ARG_SGPR, ctx->ac.i32, &ctx->gs2vs_offset);
1091 add_arg(&args, ARG_SGPR, ctx->ac.i32, &ctx->gs_wave_id);
1092 add_arg(&args, ARG_VGPR, ctx->ac.i32,
1093 &ctx->gs_vtx_offset[0]);
1094 add_arg(&args, ARG_VGPR, ctx->ac.i32,
1095 &ctx->gs_vtx_offset[1]);
1096 add_arg(&args, ARG_VGPR, ctx->ac.i32,
1097 &ctx->abi.gs_prim_id);
1098 add_arg(&args, ARG_VGPR, ctx->ac.i32,
1099 &ctx->gs_vtx_offset[2]);
1100 add_arg(&args, ARG_VGPR, ctx->ac.i32,
1101 &ctx->gs_vtx_offset[3]);
1102 add_arg(&args, ARG_VGPR, ctx->ac.i32,
1103 &ctx->gs_vtx_offset[4]);
1104 add_arg(&args, ARG_VGPR, ctx->ac.i32,
1105 &ctx->gs_vtx_offset[5]);
1106 add_arg(&args, ARG_VGPR, ctx->ac.i32,
1107 &ctx->abi.gs_invocation_id);
1108 }
1109 break;
1110 case MESA_SHADER_FRAGMENT:
1111 declare_global_input_sgprs(ctx, &user_sgpr_info, &args,
1112 &desc_sets);
1113
1114 add_arg(&args, ARG_SGPR, ctx->ac.i32, &ctx->abi.prim_mask);
1115 add_arg(&args, ARG_VGPR, ctx->ac.v2i32, &ctx->abi.persp_sample);
1116 add_arg(&args, ARG_VGPR, ctx->ac.v2i32, &ctx->abi.persp_center);
1117 add_arg(&args, ARG_VGPR, ctx->ac.v2i32, &ctx->abi.persp_centroid);
1118 add_arg(&args, ARG_VGPR, ctx->ac.v3i32, NULL); /* persp pull model */
1119 add_arg(&args, ARG_VGPR, ctx->ac.v2i32, &ctx->abi.linear_sample);
1120 add_arg(&args, ARG_VGPR, ctx->ac.v2i32, &ctx->abi.linear_center);
1121 add_arg(&args, ARG_VGPR, ctx->ac.v2i32, &ctx->abi.linear_centroid);
1122 add_arg(&args, ARG_VGPR, ctx->ac.f32, NULL); /* line stipple tex */
1123 add_arg(&args, ARG_VGPR, ctx->ac.f32, &ctx->abi.frag_pos[0]);
1124 add_arg(&args, ARG_VGPR, ctx->ac.f32, &ctx->abi.frag_pos[1]);
1125 add_arg(&args, ARG_VGPR, ctx->ac.f32, &ctx->abi.frag_pos[2]);
1126 add_arg(&args, ARG_VGPR, ctx->ac.f32, &ctx->abi.frag_pos[3]);
1127 add_arg(&args, ARG_VGPR, ctx->ac.i32, &ctx->abi.front_face);
1128 add_arg(&args, ARG_VGPR, ctx->ac.i32, &ctx->abi.ancillary);
1129 add_arg(&args, ARG_VGPR, ctx->ac.i32, &ctx->abi.sample_coverage);
1130 add_arg(&args, ARG_VGPR, ctx->ac.i32, NULL); /* fixed pt */
1131 break;
1132 default:
1133 unreachable("Shader stage not implemented");
1134 }
1135
1136 ctx->main_function = create_llvm_function(
1137 ctx->context, ctx->ac.module, ctx->ac.builder, NULL, 0, &args,
1138 ctx->max_workgroup_size, ctx->options);
1139 set_llvm_calling_convention(ctx->main_function, stage);
1140
1141
1142 ctx->shader_info->num_input_vgprs = 0;
1143 ctx->shader_info->num_input_sgprs = ctx->options->supports_spill ? 2 : 0;
1144
1145 ctx->shader_info->num_input_sgprs += args.num_sgprs_used;
1146
1147 if (ctx->stage != MESA_SHADER_FRAGMENT)
1148 ctx->shader_info->num_input_vgprs = args.num_vgprs_used;
1149
1150 assign_arguments(ctx->main_function, &args);
1151
1152 user_sgpr_idx = 0;
1153
1154 if (ctx->options->supports_spill || user_sgpr_info.need_ring_offsets) {
1155 set_loc_shader_ptr(ctx, AC_UD_SCRATCH_RING_OFFSETS,
1156 &user_sgpr_idx);
1157 if (ctx->options->supports_spill) {
1158 ctx->ring_offsets = ac_build_intrinsic(&ctx->ac, "llvm.amdgcn.implicit.buffer.ptr",
1159 LLVMPointerType(ctx->ac.i8, AC_ADDR_SPACE_CONST),
1160 NULL, 0, AC_FUNC_ATTR_READNONE);
1161 ctx->ring_offsets = LLVMBuildBitCast(ctx->ac.builder, ctx->ring_offsets,
1162 ac_array_in_const_addr_space(ctx->ac.v4i32), "");
1163 }
1164 }
1165
1166 /* For merged shaders the user SGPRs start at 8, with 8 system SGPRs in front (including
1167 * the rw_buffers at s0/s1. With user SGPR0 = s8, lets restart the count from 0 */
1168 if (has_previous_stage)
1169 user_sgpr_idx = 0;
1170
1171 set_global_input_locs(ctx, &user_sgpr_info, desc_sets, &user_sgpr_idx);
1172
1173 switch (stage) {
1174 case MESA_SHADER_COMPUTE:
1175 if (ctx->shader_info->cs.uses_grid_size) {
1176 set_loc_shader(ctx, AC_UD_CS_GRID_SIZE,
1177 &user_sgpr_idx, 3);
1178 }
1179 break;
1180 case MESA_SHADER_VERTEX:
1181 set_vs_specific_input_locs(ctx, stage, has_previous_stage,
1182 previous_stage, &user_sgpr_idx);
1183 if (ctx->abi.view_index)
1184 set_loc_shader(ctx, AC_UD_VIEW_INDEX, &user_sgpr_idx, 1);
1185 break;
1186 case MESA_SHADER_TESS_CTRL:
1187 set_vs_specific_input_locs(ctx, stage, has_previous_stage,
1188 previous_stage, &user_sgpr_idx);
1189 if (ctx->abi.view_index)
1190 set_loc_shader(ctx, AC_UD_VIEW_INDEX, &user_sgpr_idx, 1);
1191 break;
1192 case MESA_SHADER_TESS_EVAL:
1193 if (ctx->abi.view_index)
1194 set_loc_shader(ctx, AC_UD_VIEW_INDEX, &user_sgpr_idx, 1);
1195 break;
1196 case MESA_SHADER_GEOMETRY:
1197 if (has_previous_stage) {
1198 if (previous_stage == MESA_SHADER_VERTEX)
1199 set_vs_specific_input_locs(ctx, stage,
1200 has_previous_stage,
1201 previous_stage,
1202 &user_sgpr_idx);
1203 }
1204 if (ctx->abi.view_index)
1205 set_loc_shader(ctx, AC_UD_VIEW_INDEX, &user_sgpr_idx, 1);
1206 break;
1207 case MESA_SHADER_FRAGMENT:
1208 break;
1209 default:
1210 unreachable("Shader stage not implemented");
1211 }
1212
1213 if (stage == MESA_SHADER_TESS_CTRL ||
1214 (stage == MESA_SHADER_VERTEX && ctx->options->key.vs_common_out.as_ls) ||
1215 /* GFX9 has the ESGS ring buffer in LDS. */
1216 (stage == MESA_SHADER_GEOMETRY && has_previous_stage)) {
1217 ac_declare_lds_as_pointer(&ctx->ac);
1218 }
1219
1220 ctx->shader_info->num_user_sgprs = user_sgpr_idx;
1221 }
1222
1223
1224 static LLVMValueRef
1225 radv_load_resource(struct ac_shader_abi *abi, LLVMValueRef index,
1226 unsigned desc_set, unsigned binding)
1227 {
1228 struct radv_shader_context *ctx = radv_shader_context_from_abi(abi);
1229 LLVMValueRef desc_ptr = ctx->descriptor_sets[desc_set];
1230 struct radv_pipeline_layout *pipeline_layout = ctx->options->layout;
1231 struct radv_descriptor_set_layout *layout = pipeline_layout->set[desc_set].layout;
1232 unsigned base_offset = layout->binding[binding].offset;
1233 LLVMValueRef offset, stride;
1234
1235 if (layout->binding[binding].type == VK_DESCRIPTOR_TYPE_UNIFORM_BUFFER_DYNAMIC ||
1236 layout->binding[binding].type == VK_DESCRIPTOR_TYPE_STORAGE_BUFFER_DYNAMIC) {
1237 unsigned idx = pipeline_layout->set[desc_set].dynamic_offset_start +
1238 layout->binding[binding].dynamic_offset_offset;
1239 desc_ptr = ctx->abi.push_constants;
1240 base_offset = pipeline_layout->push_constant_size + 16 * idx;
1241 stride = LLVMConstInt(ctx->ac.i32, 16, false);
1242 } else
1243 stride = LLVMConstInt(ctx->ac.i32, layout->binding[binding].size, false);
1244
1245 offset = LLVMConstInt(ctx->ac.i32, base_offset, false);
1246
1247 if (layout->binding[binding].type != VK_DESCRIPTOR_TYPE_INLINE_UNIFORM_BLOCK_EXT) {
1248 offset = ac_build_imad(&ctx->ac, index, stride, offset);
1249 }
1250
1251 desc_ptr = LLVMBuildGEP(ctx->ac.builder, desc_ptr, &offset, 1, "");
1252 desc_ptr = ac_cast_ptr(&ctx->ac, desc_ptr, ctx->ac.v4i32);
1253 LLVMSetMetadata(desc_ptr, ctx->ac.uniform_md_kind, ctx->ac.empty_md);
1254
1255 if (layout->binding[binding].type == VK_DESCRIPTOR_TYPE_INLINE_UNIFORM_BLOCK_EXT) {
1256 uint32_t desc_type = S_008F0C_DST_SEL_X(V_008F0C_SQ_SEL_X) |
1257 S_008F0C_DST_SEL_Y(V_008F0C_SQ_SEL_Y) |
1258 S_008F0C_DST_SEL_Z(V_008F0C_SQ_SEL_Z) |
1259 S_008F0C_DST_SEL_W(V_008F0C_SQ_SEL_W);
1260
1261 if (ctx->ac.chip_class >= GFX10) {
1262 desc_type |= S_008F0C_FORMAT(V_008F0C_IMG_FORMAT_32_FLOAT) |
1263 S_008F0C_OOB_SELECT(3) |
1264 S_008F0C_RESOURCE_LEVEL(1);
1265 } else {
1266 desc_type |= S_008F0C_NUM_FORMAT(V_008F0C_BUF_NUM_FORMAT_FLOAT) |
1267 S_008F0C_DATA_FORMAT(V_008F0C_BUF_DATA_FORMAT_32);
1268 }
1269
1270 LLVMValueRef desc_components[4] = {
1271 LLVMBuildPtrToInt(ctx->ac.builder, desc_ptr, ctx->ac.intptr, ""),
1272 LLVMConstInt(ctx->ac.i32, S_008F04_BASE_ADDRESS_HI(ctx->options->address32_hi), false),
1273 /* High limit to support variable sizes. */
1274 LLVMConstInt(ctx->ac.i32, 0xffffffff, false),
1275 LLVMConstInt(ctx->ac.i32, desc_type, false),
1276 };
1277
1278 return ac_build_gather_values(&ctx->ac, desc_components, 4);
1279 }
1280
1281 return desc_ptr;
1282 }
1283
1284
1285 /* The offchip buffer layout for TCS->TES is
1286 *
1287 * - attribute 0 of patch 0 vertex 0
1288 * - attribute 0 of patch 0 vertex 1
1289 * - attribute 0 of patch 0 vertex 2
1290 * ...
1291 * - attribute 0 of patch 1 vertex 0
1292 * - attribute 0 of patch 1 vertex 1
1293 * ...
1294 * - attribute 1 of patch 0 vertex 0
1295 * - attribute 1 of patch 0 vertex 1
1296 * ...
1297 * - per patch attribute 0 of patch 0
1298 * - per patch attribute 0 of patch 1
1299 * ...
1300 *
1301 * Note that every attribute has 4 components.
1302 */
1303 static LLVMValueRef get_non_vertex_index_offset(struct radv_shader_context *ctx)
1304 {
1305 uint32_t num_patches = ctx->tcs_num_patches;
1306 uint32_t num_tcs_outputs;
1307 if (ctx->stage == MESA_SHADER_TESS_CTRL)
1308 num_tcs_outputs = util_last_bit64(ctx->shader_info->tcs.outputs_written);
1309 else
1310 num_tcs_outputs = ctx->options->key.tes.tcs_num_outputs;
1311
1312 uint32_t output_vertex_size = num_tcs_outputs * 16;
1313 uint32_t pervertex_output_patch_size = ctx->shader->info.tess.tcs_vertices_out * output_vertex_size;
1314
1315 return LLVMConstInt(ctx->ac.i32, pervertex_output_patch_size * num_patches, false);
1316 }
1317
1318 static LLVMValueRef calc_param_stride(struct radv_shader_context *ctx,
1319 LLVMValueRef vertex_index)
1320 {
1321 LLVMValueRef param_stride;
1322 if (vertex_index)
1323 param_stride = LLVMConstInt(ctx->ac.i32, ctx->shader->info.tess.tcs_vertices_out * ctx->tcs_num_patches, false);
1324 else
1325 param_stride = LLVMConstInt(ctx->ac.i32, ctx->tcs_num_patches, false);
1326 return param_stride;
1327 }
1328
1329 static LLVMValueRef get_tcs_tes_buffer_address(struct radv_shader_context *ctx,
1330 LLVMValueRef vertex_index,
1331 LLVMValueRef param_index)
1332 {
1333 LLVMValueRef base_addr;
1334 LLVMValueRef param_stride, constant16;
1335 LLVMValueRef rel_patch_id = get_rel_patch_id(ctx);
1336 LLVMValueRef vertices_per_patch = LLVMConstInt(ctx->ac.i32, ctx->shader->info.tess.tcs_vertices_out, false);
1337 constant16 = LLVMConstInt(ctx->ac.i32, 16, false);
1338 param_stride = calc_param_stride(ctx, vertex_index);
1339 if (vertex_index) {
1340 base_addr = ac_build_imad(&ctx->ac, rel_patch_id,
1341 vertices_per_patch, vertex_index);
1342 } else {
1343 base_addr = rel_patch_id;
1344 }
1345
1346 base_addr = LLVMBuildAdd(ctx->ac.builder, base_addr,
1347 LLVMBuildMul(ctx->ac.builder, param_index,
1348 param_stride, ""), "");
1349
1350 base_addr = LLVMBuildMul(ctx->ac.builder, base_addr, constant16, "");
1351
1352 if (!vertex_index) {
1353 LLVMValueRef patch_data_offset = get_non_vertex_index_offset(ctx);
1354
1355 base_addr = LLVMBuildAdd(ctx->ac.builder, base_addr,
1356 patch_data_offset, "");
1357 }
1358 return base_addr;
1359 }
1360
1361 static LLVMValueRef get_tcs_tes_buffer_address_params(struct radv_shader_context *ctx,
1362 unsigned param,
1363 unsigned const_index,
1364 bool is_compact,
1365 LLVMValueRef vertex_index,
1366 LLVMValueRef indir_index)
1367 {
1368 LLVMValueRef param_index;
1369
1370 if (indir_index)
1371 param_index = LLVMBuildAdd(ctx->ac.builder, LLVMConstInt(ctx->ac.i32, param, false),
1372 indir_index, "");
1373 else {
1374 if (const_index && !is_compact)
1375 param += const_index;
1376 param_index = LLVMConstInt(ctx->ac.i32, param, false);
1377 }
1378 return get_tcs_tes_buffer_address(ctx, vertex_index, param_index);
1379 }
1380
1381 static LLVMValueRef
1382 get_dw_address(struct radv_shader_context *ctx,
1383 LLVMValueRef dw_addr,
1384 unsigned param,
1385 unsigned const_index,
1386 bool compact_const_index,
1387 LLVMValueRef vertex_index,
1388 LLVMValueRef stride,
1389 LLVMValueRef indir_index)
1390
1391 {
1392
1393 if (vertex_index) {
1394 dw_addr = LLVMBuildAdd(ctx->ac.builder, dw_addr,
1395 LLVMBuildMul(ctx->ac.builder,
1396 vertex_index,
1397 stride, ""), "");
1398 }
1399
1400 if (indir_index)
1401 dw_addr = LLVMBuildAdd(ctx->ac.builder, dw_addr,
1402 LLVMBuildMul(ctx->ac.builder, indir_index,
1403 LLVMConstInt(ctx->ac.i32, 4, false), ""), "");
1404 else if (const_index && !compact_const_index)
1405 dw_addr = LLVMBuildAdd(ctx->ac.builder, dw_addr,
1406 LLVMConstInt(ctx->ac.i32, const_index * 4, false), "");
1407
1408 dw_addr = LLVMBuildAdd(ctx->ac.builder, dw_addr,
1409 LLVMConstInt(ctx->ac.i32, param * 4, false), "");
1410
1411 if (const_index && compact_const_index)
1412 dw_addr = LLVMBuildAdd(ctx->ac.builder, dw_addr,
1413 LLVMConstInt(ctx->ac.i32, const_index, false), "");
1414 return dw_addr;
1415 }
1416
1417 static LLVMValueRef
1418 load_tcs_varyings(struct ac_shader_abi *abi,
1419 LLVMTypeRef type,
1420 LLVMValueRef vertex_index,
1421 LLVMValueRef indir_index,
1422 unsigned const_index,
1423 unsigned location,
1424 unsigned driver_location,
1425 unsigned component,
1426 unsigned num_components,
1427 bool is_patch,
1428 bool is_compact,
1429 bool load_input)
1430 {
1431 struct radv_shader_context *ctx = radv_shader_context_from_abi(abi);
1432 LLVMValueRef dw_addr, stride;
1433 LLVMValueRef value[4], result;
1434 unsigned param = shader_io_get_unique_index(location);
1435
1436 if (load_input) {
1437 uint32_t input_vertex_size = (ctx->tcs_num_inputs * 16) / 4;
1438 stride = LLVMConstInt(ctx->ac.i32, input_vertex_size, false);
1439 dw_addr = get_tcs_in_current_patch_offset(ctx);
1440 } else {
1441 if (!is_patch) {
1442 stride = get_tcs_out_vertex_stride(ctx);
1443 dw_addr = get_tcs_out_current_patch_offset(ctx);
1444 } else {
1445 dw_addr = get_tcs_out_current_patch_data_offset(ctx);
1446 stride = NULL;
1447 }
1448 }
1449
1450 dw_addr = get_dw_address(ctx, dw_addr, param, const_index, is_compact, vertex_index, stride,
1451 indir_index);
1452
1453 for (unsigned i = 0; i < num_components + component; i++) {
1454 value[i] = ac_lds_load(&ctx->ac, dw_addr);
1455 dw_addr = LLVMBuildAdd(ctx->ac.builder, dw_addr,
1456 ctx->ac.i32_1, "");
1457 }
1458 result = ac_build_varying_gather_values(&ctx->ac, value, num_components, component);
1459 return result;
1460 }
1461
1462 static void
1463 store_tcs_output(struct ac_shader_abi *abi,
1464 const nir_variable *var,
1465 LLVMValueRef vertex_index,
1466 LLVMValueRef param_index,
1467 unsigned const_index,
1468 LLVMValueRef src,
1469 unsigned writemask)
1470 {
1471 struct radv_shader_context *ctx = radv_shader_context_from_abi(abi);
1472 const unsigned location = var->data.location;
1473 unsigned component = var->data.location_frac;
1474 const bool is_patch = var->data.patch;
1475 const bool is_compact = var->data.compact;
1476 LLVMValueRef dw_addr;
1477 LLVMValueRef stride = NULL;
1478 LLVMValueRef buf_addr = NULL;
1479 unsigned param;
1480 bool store_lds = true;
1481
1482 if (is_patch) {
1483 if (!(ctx->shader->info.patch_outputs_read & (1U << (location - VARYING_SLOT_PATCH0))))
1484 store_lds = false;
1485 } else {
1486 if (!(ctx->shader->info.outputs_read & (1ULL << location)))
1487 store_lds = false;
1488 }
1489
1490 param = shader_io_get_unique_index(location);
1491 if ((location == VARYING_SLOT_CLIP_DIST0 || location == VARYING_SLOT_CLIP_DIST1) && is_compact) {
1492 const_index += component;
1493 component = 0;
1494
1495 if (const_index >= 4) {
1496 const_index -= 4;
1497 param++;
1498 }
1499 }
1500
1501 if (!is_patch) {
1502 stride = get_tcs_out_vertex_stride(ctx);
1503 dw_addr = get_tcs_out_current_patch_offset(ctx);
1504 } else {
1505 dw_addr = get_tcs_out_current_patch_data_offset(ctx);
1506 }
1507
1508 dw_addr = get_dw_address(ctx, dw_addr, param, const_index, is_compact, vertex_index, stride,
1509 param_index);
1510 buf_addr = get_tcs_tes_buffer_address_params(ctx, param, const_index, is_compact,
1511 vertex_index, param_index);
1512
1513 bool is_tess_factor = false;
1514 if (location == VARYING_SLOT_TESS_LEVEL_INNER ||
1515 location == VARYING_SLOT_TESS_LEVEL_OUTER)
1516 is_tess_factor = true;
1517
1518 unsigned base = is_compact ? const_index : 0;
1519 for (unsigned chan = 0; chan < 8; chan++) {
1520 if (!(writemask & (1 << chan)))
1521 continue;
1522 LLVMValueRef value = ac_llvm_extract_elem(&ctx->ac, src, chan - component);
1523 value = ac_to_integer(&ctx->ac, value);
1524 value = LLVMBuildZExtOrBitCast(ctx->ac.builder, value, ctx->ac.i32, "");
1525
1526 if (store_lds || is_tess_factor) {
1527 LLVMValueRef dw_addr_chan =
1528 LLVMBuildAdd(ctx->ac.builder, dw_addr,
1529 LLVMConstInt(ctx->ac.i32, chan, false), "");
1530 ac_lds_store(&ctx->ac, dw_addr_chan, value);
1531 }
1532
1533 if (!is_tess_factor && writemask != 0xF)
1534 ac_build_buffer_store_dword(&ctx->ac, ctx->hs_ring_tess_offchip, value, 1,
1535 buf_addr, ctx->oc_lds,
1536 4 * (base + chan), ac_glc, false);
1537 }
1538
1539 if (writemask == 0xF) {
1540 ac_build_buffer_store_dword(&ctx->ac, ctx->hs_ring_tess_offchip, src, 4,
1541 buf_addr, ctx->oc_lds,
1542 (base * 4), ac_glc, false);
1543 }
1544 }
1545
1546 static LLVMValueRef
1547 load_tes_input(struct ac_shader_abi *abi,
1548 LLVMTypeRef type,
1549 LLVMValueRef vertex_index,
1550 LLVMValueRef param_index,
1551 unsigned const_index,
1552 unsigned location,
1553 unsigned driver_location,
1554 unsigned component,
1555 unsigned num_components,
1556 bool is_patch,
1557 bool is_compact,
1558 bool load_input)
1559 {
1560 struct radv_shader_context *ctx = radv_shader_context_from_abi(abi);
1561 LLVMValueRef buf_addr;
1562 LLVMValueRef result;
1563 unsigned param = shader_io_get_unique_index(location);
1564
1565 if ((location == VARYING_SLOT_CLIP_DIST0 || location == VARYING_SLOT_CLIP_DIST1) && is_compact) {
1566 const_index += component;
1567 component = 0;
1568 if (const_index >= 4) {
1569 const_index -= 4;
1570 param++;
1571 }
1572 }
1573
1574 buf_addr = get_tcs_tes_buffer_address_params(ctx, param, const_index,
1575 is_compact, vertex_index, param_index);
1576
1577 LLVMValueRef comp_offset = LLVMConstInt(ctx->ac.i32, component * 4, false);
1578 buf_addr = LLVMBuildAdd(ctx->ac.builder, buf_addr, comp_offset, "");
1579
1580 result = ac_build_buffer_load(&ctx->ac, ctx->hs_ring_tess_offchip, num_components, NULL,
1581 buf_addr, ctx->oc_lds, is_compact ? (4 * const_index) : 0, ac_glc, true, false);
1582 result = ac_trim_vector(&ctx->ac, result, num_components);
1583 return result;
1584 }
1585
1586 static LLVMValueRef
1587 radv_emit_fetch_64bit(struct radv_shader_context *ctx,
1588 LLVMTypeRef type, LLVMValueRef a, LLVMValueRef b)
1589 {
1590 LLVMValueRef values[2] = {
1591 ac_to_integer(&ctx->ac, a),
1592 ac_to_integer(&ctx->ac, b),
1593 };
1594 LLVMValueRef result = ac_build_gather_values(&ctx->ac, values, 2);
1595 return LLVMBuildBitCast(ctx->ac.builder, result, type, "");
1596 }
1597
1598 static LLVMValueRef
1599 load_gs_input(struct ac_shader_abi *abi,
1600 unsigned location,
1601 unsigned driver_location,
1602 unsigned component,
1603 unsigned num_components,
1604 unsigned vertex_index,
1605 unsigned const_index,
1606 LLVMTypeRef type)
1607 {
1608 struct radv_shader_context *ctx = radv_shader_context_from_abi(abi);
1609 LLVMValueRef vtx_offset;
1610 unsigned param, vtx_offset_param;
1611 LLVMValueRef value[4], result;
1612
1613 vtx_offset_param = vertex_index;
1614 assert(vtx_offset_param < 6);
1615 vtx_offset = LLVMBuildMul(ctx->ac.builder, ctx->gs_vtx_offset[vtx_offset_param],
1616 LLVMConstInt(ctx->ac.i32, 4, false), "");
1617
1618 param = shader_io_get_unique_index(location);
1619
1620 for (unsigned i = component; i < num_components + component; i++) {
1621 if (ctx->ac.chip_class >= GFX9) {
1622 LLVMValueRef dw_addr = ctx->gs_vtx_offset[vtx_offset_param];
1623 dw_addr = LLVMBuildAdd(ctx->ac.builder, dw_addr,
1624 LLVMConstInt(ctx->ac.i32, param * 4 + i + const_index, 0), "");
1625 value[i] = ac_lds_load(&ctx->ac, dw_addr);
1626
1627 if (ac_get_type_size(type) == 8) {
1628 dw_addr = LLVMBuildAdd(ctx->ac.builder, dw_addr,
1629 LLVMConstInt(ctx->ac.i32, param * 4 + i + const_index + 1, 0), "");
1630 LLVMValueRef tmp = ac_lds_load(&ctx->ac, dw_addr);
1631
1632 value[i] = radv_emit_fetch_64bit(ctx, type, value[i], tmp);
1633 }
1634 } else {
1635 LLVMValueRef soffset =
1636 LLVMConstInt(ctx->ac.i32,
1637 (param * 4 + i + const_index) * 256,
1638 false);
1639
1640 value[i] = ac_build_buffer_load(&ctx->ac,
1641 ctx->esgs_ring, 1,
1642 ctx->ac.i32_0,
1643 vtx_offset, soffset,
1644 0, ac_glc, true, false);
1645
1646 if (ac_get_type_size(type) == 8) {
1647 soffset = LLVMConstInt(ctx->ac.i32,
1648 (param * 4 + i + const_index + 1) * 256,
1649 false);
1650
1651 LLVMValueRef tmp =
1652 ac_build_buffer_load(&ctx->ac,
1653 ctx->esgs_ring, 1,
1654 ctx->ac.i32_0,
1655 vtx_offset, soffset,
1656 0, ac_glc, true, false);
1657
1658 value[i] = radv_emit_fetch_64bit(ctx, type, value[i], tmp);
1659 }
1660 }
1661
1662 if (ac_get_type_size(type) == 2) {
1663 value[i] = LLVMBuildBitCast(ctx->ac.builder, value[i], ctx->ac.i32, "");
1664 value[i] = LLVMBuildTrunc(ctx->ac.builder, value[i], ctx->ac.i16, "");
1665 }
1666 value[i] = LLVMBuildBitCast(ctx->ac.builder, value[i], type, "");
1667 }
1668 result = ac_build_varying_gather_values(&ctx->ac, value, num_components, component);
1669 result = ac_to_integer(&ctx->ac, result);
1670 return result;
1671 }
1672
1673
1674 static void radv_emit_kill(struct ac_shader_abi *abi, LLVMValueRef visible)
1675 {
1676 struct radv_shader_context *ctx = radv_shader_context_from_abi(abi);
1677 ac_build_kill_if_false(&ctx->ac, visible);
1678 }
1679
1680 static uint32_t
1681 radv_get_sample_pos_offset(uint32_t num_samples)
1682 {
1683 uint32_t sample_pos_offset = 0;
1684
1685 switch (num_samples) {
1686 case 2:
1687 sample_pos_offset = 1;
1688 break;
1689 case 4:
1690 sample_pos_offset = 3;
1691 break;
1692 case 8:
1693 sample_pos_offset = 7;
1694 break;
1695 default:
1696 break;
1697 }
1698 return sample_pos_offset;
1699 }
1700
1701 static LLVMValueRef load_sample_position(struct ac_shader_abi *abi,
1702 LLVMValueRef sample_id)
1703 {
1704 struct radv_shader_context *ctx = radv_shader_context_from_abi(abi);
1705
1706 LLVMValueRef result;
1707 LLVMValueRef index = LLVMConstInt(ctx->ac.i32, RING_PS_SAMPLE_POSITIONS, false);
1708 LLVMValueRef ptr = LLVMBuildGEP(ctx->ac.builder, ctx->ring_offsets, &index, 1, "");
1709
1710 ptr = LLVMBuildBitCast(ctx->ac.builder, ptr,
1711 ac_array_in_const_addr_space(ctx->ac.v2f32), "");
1712
1713 uint32_t sample_pos_offset =
1714 radv_get_sample_pos_offset(ctx->options->key.fs.num_samples);
1715
1716 sample_id =
1717 LLVMBuildAdd(ctx->ac.builder, sample_id,
1718 LLVMConstInt(ctx->ac.i32, sample_pos_offset, false), "");
1719 result = ac_build_load_invariant(&ctx->ac, ptr, sample_id);
1720
1721 return result;
1722 }
1723
1724
1725 static LLVMValueRef load_sample_mask_in(struct ac_shader_abi *abi)
1726 {
1727 struct radv_shader_context *ctx = radv_shader_context_from_abi(abi);
1728 uint8_t log2_ps_iter_samples;
1729
1730 if (ctx->shader_info->ps.force_persample) {
1731 log2_ps_iter_samples =
1732 util_logbase2(ctx->options->key.fs.num_samples);
1733 } else {
1734 log2_ps_iter_samples = ctx->options->key.fs.log2_ps_iter_samples;
1735 }
1736
1737 /* The bit pattern matches that used by fixed function fragment
1738 * processing. */
1739 static const uint16_t ps_iter_masks[] = {
1740 0xffff, /* not used */
1741 0x5555,
1742 0x1111,
1743 0x0101,
1744 0x0001,
1745 };
1746 assert(log2_ps_iter_samples < ARRAY_SIZE(ps_iter_masks));
1747
1748 uint32_t ps_iter_mask = ps_iter_masks[log2_ps_iter_samples];
1749
1750 LLVMValueRef result, sample_id;
1751 sample_id = ac_unpack_param(&ctx->ac, abi->ancillary, 8, 4);
1752 sample_id = LLVMBuildShl(ctx->ac.builder, LLVMConstInt(ctx->ac.i32, ps_iter_mask, false), sample_id, "");
1753 result = LLVMBuildAnd(ctx->ac.builder, sample_id, abi->sample_coverage, "");
1754 return result;
1755 }
1756
1757
1758 static void gfx10_ngg_gs_emit_vertex(struct radv_shader_context *ctx,
1759 unsigned stream,
1760 LLVMValueRef *addrs);
1761
1762 static void
1763 visit_emit_vertex(struct ac_shader_abi *abi, unsigned stream, LLVMValueRef *addrs)
1764 {
1765 LLVMValueRef gs_next_vertex;
1766 LLVMValueRef can_emit;
1767 unsigned offset = 0;
1768 struct radv_shader_context *ctx = radv_shader_context_from_abi(abi);
1769
1770 if (ctx->options->key.vs_common_out.as_ngg) {
1771 gfx10_ngg_gs_emit_vertex(ctx, stream, addrs);
1772 return;
1773 }
1774
1775 /* Write vertex attribute values to GSVS ring */
1776 gs_next_vertex = LLVMBuildLoad(ctx->ac.builder,
1777 ctx->gs_next_vertex[stream],
1778 "");
1779
1780 /* If this thread has already emitted the declared maximum number of
1781 * vertices, don't emit any more: excessive vertex emissions are not
1782 * supposed to have any effect.
1783 */
1784 can_emit = LLVMBuildICmp(ctx->ac.builder, LLVMIntULT, gs_next_vertex,
1785 LLVMConstInt(ctx->ac.i32, ctx->shader->info.gs.vertices_out, false), "");
1786
1787 bool use_kill = !ctx->shader_info->gs.writes_memory;
1788 if (use_kill)
1789 ac_build_kill_if_false(&ctx->ac, can_emit);
1790 else
1791 ac_build_ifcc(&ctx->ac, can_emit, 6505);
1792
1793 for (unsigned i = 0; i < AC_LLVM_MAX_OUTPUTS; ++i) {
1794 unsigned output_usage_mask =
1795 ctx->shader_info->gs.output_usage_mask[i];
1796 uint8_t output_stream =
1797 ctx->shader_info->gs.output_streams[i];
1798 LLVMValueRef *out_ptr = &addrs[i * 4];
1799 int length = util_last_bit(output_usage_mask);
1800
1801 if (!(ctx->output_mask & (1ull << i)) ||
1802 output_stream != stream)
1803 continue;
1804
1805 for (unsigned j = 0; j < length; j++) {
1806 if (!(output_usage_mask & (1 << j)))
1807 continue;
1808
1809 LLVMValueRef out_val = LLVMBuildLoad(ctx->ac.builder,
1810 out_ptr[j], "");
1811 LLVMValueRef voffset =
1812 LLVMConstInt(ctx->ac.i32, offset *
1813 ctx->shader->info.gs.vertices_out, false);
1814
1815 offset++;
1816
1817 voffset = LLVMBuildAdd(ctx->ac.builder, voffset, gs_next_vertex, "");
1818 voffset = LLVMBuildMul(ctx->ac.builder, voffset, LLVMConstInt(ctx->ac.i32, 4, false), "");
1819
1820 out_val = ac_to_integer(&ctx->ac, out_val);
1821 out_val = LLVMBuildZExtOrBitCast(ctx->ac.builder, out_val, ctx->ac.i32, "");
1822
1823 ac_build_buffer_store_dword(&ctx->ac,
1824 ctx->gsvs_ring[stream],
1825 out_val, 1,
1826 voffset, ctx->gs2vs_offset, 0,
1827 ac_glc | ac_slc, true);
1828 }
1829 }
1830
1831 gs_next_vertex = LLVMBuildAdd(ctx->ac.builder, gs_next_vertex,
1832 ctx->ac.i32_1, "");
1833 LLVMBuildStore(ctx->ac.builder, gs_next_vertex, ctx->gs_next_vertex[stream]);
1834
1835 ac_build_sendmsg(&ctx->ac,
1836 AC_SENDMSG_GS_OP_EMIT | AC_SENDMSG_GS | (stream << 8),
1837 ctx->gs_wave_id);
1838
1839 if (!use_kill)
1840 ac_build_endif(&ctx->ac, 6505);
1841 }
1842
1843 static void
1844 visit_end_primitive(struct ac_shader_abi *abi, unsigned stream)
1845 {
1846 struct radv_shader_context *ctx = radv_shader_context_from_abi(abi);
1847
1848 if (ctx->options->key.vs_common_out.as_ngg) {
1849 LLVMBuildStore(ctx->ac.builder, ctx->ac.i32_0, ctx->gs_curprim_verts[stream]);
1850 return;
1851 }
1852
1853 ac_build_sendmsg(&ctx->ac, AC_SENDMSG_GS_OP_CUT | AC_SENDMSG_GS | (stream << 8), ctx->gs_wave_id);
1854 }
1855
1856 static LLVMValueRef
1857 load_tess_coord(struct ac_shader_abi *abi)
1858 {
1859 struct radv_shader_context *ctx = radv_shader_context_from_abi(abi);
1860
1861 LLVMValueRef coord[4] = {
1862 ctx->tes_u,
1863 ctx->tes_v,
1864 ctx->ac.f32_0,
1865 ctx->ac.f32_0,
1866 };
1867
1868 if (ctx->shader->info.tess.primitive_mode == GL_TRIANGLES)
1869 coord[2] = LLVMBuildFSub(ctx->ac.builder, ctx->ac.f32_1,
1870 LLVMBuildFAdd(ctx->ac.builder, coord[0], coord[1], ""), "");
1871
1872 return ac_build_gather_values(&ctx->ac, coord, 3);
1873 }
1874
1875 static LLVMValueRef
1876 load_patch_vertices_in(struct ac_shader_abi *abi)
1877 {
1878 struct radv_shader_context *ctx = radv_shader_context_from_abi(abi);
1879 return LLVMConstInt(ctx->ac.i32, ctx->options->key.tcs.input_vertices, false);
1880 }
1881
1882
1883 static LLVMValueRef radv_load_base_vertex(struct ac_shader_abi *abi)
1884 {
1885 return abi->base_vertex;
1886 }
1887
1888 static LLVMValueRef radv_load_ssbo(struct ac_shader_abi *abi,
1889 LLVMValueRef buffer_ptr, bool write)
1890 {
1891 struct radv_shader_context *ctx = radv_shader_context_from_abi(abi);
1892 LLVMValueRef result;
1893
1894 LLVMSetMetadata(buffer_ptr, ctx->ac.uniform_md_kind, ctx->ac.empty_md);
1895
1896 result = LLVMBuildLoad(ctx->ac.builder, buffer_ptr, "");
1897 LLVMSetMetadata(result, ctx->ac.invariant_load_md_kind, ctx->ac.empty_md);
1898
1899 return result;
1900 }
1901
1902 static LLVMValueRef radv_load_ubo(struct ac_shader_abi *abi, LLVMValueRef buffer_ptr)
1903 {
1904 struct radv_shader_context *ctx = radv_shader_context_from_abi(abi);
1905 LLVMValueRef result;
1906
1907 if (LLVMGetTypeKind(LLVMTypeOf(buffer_ptr)) != LLVMPointerTypeKind) {
1908 /* Do not load the descriptor for inlined uniform blocks. */
1909 return buffer_ptr;
1910 }
1911
1912 LLVMSetMetadata(buffer_ptr, ctx->ac.uniform_md_kind, ctx->ac.empty_md);
1913
1914 result = LLVMBuildLoad(ctx->ac.builder, buffer_ptr, "");
1915 LLVMSetMetadata(result, ctx->ac.invariant_load_md_kind, ctx->ac.empty_md);
1916
1917 return result;
1918 }
1919
1920 static LLVMValueRef radv_get_sampler_desc(struct ac_shader_abi *abi,
1921 unsigned descriptor_set,
1922 unsigned base_index,
1923 unsigned constant_index,
1924 LLVMValueRef index,
1925 enum ac_descriptor_type desc_type,
1926 bool image, bool write,
1927 bool bindless)
1928 {
1929 struct radv_shader_context *ctx = radv_shader_context_from_abi(abi);
1930 LLVMValueRef list = ctx->descriptor_sets[descriptor_set];
1931 struct radv_descriptor_set_layout *layout = ctx->options->layout->set[descriptor_set].layout;
1932 struct radv_descriptor_set_binding_layout *binding = layout->binding + base_index;
1933 unsigned offset = binding->offset;
1934 unsigned stride = binding->size;
1935 unsigned type_size;
1936 LLVMBuilderRef builder = ctx->ac.builder;
1937 LLVMTypeRef type;
1938
1939 assert(base_index < layout->binding_count);
1940
1941 switch (desc_type) {
1942 case AC_DESC_IMAGE:
1943 type = ctx->ac.v8i32;
1944 type_size = 32;
1945 break;
1946 case AC_DESC_FMASK:
1947 type = ctx->ac.v8i32;
1948 offset += 32;
1949 type_size = 32;
1950 break;
1951 case AC_DESC_SAMPLER:
1952 type = ctx->ac.v4i32;
1953 if (binding->type == VK_DESCRIPTOR_TYPE_COMBINED_IMAGE_SAMPLER) {
1954 offset += radv_combined_image_descriptor_sampler_offset(binding);
1955 }
1956
1957 type_size = 16;
1958 break;
1959 case AC_DESC_BUFFER:
1960 type = ctx->ac.v4i32;
1961 type_size = 16;
1962 break;
1963 case AC_DESC_PLANE_0:
1964 case AC_DESC_PLANE_1:
1965 case AC_DESC_PLANE_2:
1966 type = ctx->ac.v8i32;
1967 type_size = 32;
1968 offset += 32 * (desc_type - AC_DESC_PLANE_0);
1969 break;
1970 default:
1971 unreachable("invalid desc_type\n");
1972 }
1973
1974 offset += constant_index * stride;
1975
1976 if (desc_type == AC_DESC_SAMPLER && binding->immutable_samplers_offset &&
1977 (!index || binding->immutable_samplers_equal)) {
1978 if (binding->immutable_samplers_equal)
1979 constant_index = 0;
1980
1981 const uint32_t *samplers = radv_immutable_samplers(layout, binding);
1982
1983 LLVMValueRef constants[] = {
1984 LLVMConstInt(ctx->ac.i32, samplers[constant_index * 4 + 0], 0),
1985 LLVMConstInt(ctx->ac.i32, samplers[constant_index * 4 + 1], 0),
1986 LLVMConstInt(ctx->ac.i32, samplers[constant_index * 4 + 2], 0),
1987 LLVMConstInt(ctx->ac.i32, samplers[constant_index * 4 + 3], 0),
1988 };
1989 return ac_build_gather_values(&ctx->ac, constants, 4);
1990 }
1991
1992 assert(stride % type_size == 0);
1993
1994 LLVMValueRef adjusted_index = index;
1995 if (!adjusted_index)
1996 adjusted_index = ctx->ac.i32_0;
1997
1998 adjusted_index = LLVMBuildMul(builder, adjusted_index, LLVMConstInt(ctx->ac.i32, stride / type_size, 0), "");
1999
2000 LLVMValueRef val_offset = LLVMConstInt(ctx->ac.i32, offset, 0);
2001 list = LLVMBuildGEP(builder, list, &val_offset, 1, "");
2002 list = LLVMBuildPointerCast(builder, list,
2003 ac_array_in_const32_addr_space(type), "");
2004
2005 LLVMValueRef descriptor = ac_build_load_to_sgpr(&ctx->ac, list, adjusted_index);
2006
2007 /* 3 plane formats always have same size and format for plane 1 & 2, so
2008 * use the tail from plane 1 so that we can store only the first 16 bytes
2009 * of the last plane. */
2010 if (desc_type == AC_DESC_PLANE_2) {
2011 LLVMValueRef descriptor2 = radv_get_sampler_desc(abi, descriptor_set, base_index, constant_index, index, AC_DESC_PLANE_1,image, write, bindless);
2012
2013 LLVMValueRef components[8];
2014 for (unsigned i = 0; i < 4; ++i)
2015 components[i] = ac_llvm_extract_elem(&ctx->ac, descriptor, i);
2016
2017 for (unsigned i = 4; i < 8; ++i)
2018 components[i] = ac_llvm_extract_elem(&ctx->ac, descriptor2, i);
2019 descriptor = ac_build_gather_values(&ctx->ac, components, 8);
2020 }
2021
2022 return descriptor;
2023 }
2024
2025 /* For 2_10_10_10 formats the alpha is handled as unsigned by pre-vega HW.
2026 * so we may need to fix it up. */
2027 static LLVMValueRef
2028 adjust_vertex_fetch_alpha(struct radv_shader_context *ctx,
2029 unsigned adjustment,
2030 LLVMValueRef alpha)
2031 {
2032 if (adjustment == RADV_ALPHA_ADJUST_NONE)
2033 return alpha;
2034
2035 LLVMValueRef c30 = LLVMConstInt(ctx->ac.i32, 30, 0);
2036
2037 alpha = LLVMBuildBitCast(ctx->ac.builder, alpha, ctx->ac.f32, "");
2038
2039 if (adjustment == RADV_ALPHA_ADJUST_SSCALED)
2040 alpha = LLVMBuildFPToUI(ctx->ac.builder, alpha, ctx->ac.i32, "");
2041 else
2042 alpha = ac_to_integer(&ctx->ac, alpha);
2043
2044 /* For the integer-like cases, do a natural sign extension.
2045 *
2046 * For the SNORM case, the values are 0.0, 0.333, 0.666, 1.0
2047 * and happen to contain 0, 1, 2, 3 as the two LSBs of the
2048 * exponent.
2049 */
2050 alpha = LLVMBuildShl(ctx->ac.builder, alpha,
2051 adjustment == RADV_ALPHA_ADJUST_SNORM ?
2052 LLVMConstInt(ctx->ac.i32, 7, 0) : c30, "");
2053 alpha = LLVMBuildAShr(ctx->ac.builder, alpha, c30, "");
2054
2055 /* Convert back to the right type. */
2056 if (adjustment == RADV_ALPHA_ADJUST_SNORM) {
2057 LLVMValueRef clamp;
2058 LLVMValueRef neg_one = LLVMConstReal(ctx->ac.f32, -1.0);
2059 alpha = LLVMBuildSIToFP(ctx->ac.builder, alpha, ctx->ac.f32, "");
2060 clamp = LLVMBuildFCmp(ctx->ac.builder, LLVMRealULT, alpha, neg_one, "");
2061 alpha = LLVMBuildSelect(ctx->ac.builder, clamp, neg_one, alpha, "");
2062 } else if (adjustment == RADV_ALPHA_ADJUST_SSCALED) {
2063 alpha = LLVMBuildSIToFP(ctx->ac.builder, alpha, ctx->ac.f32, "");
2064 }
2065
2066 return LLVMBuildBitCast(ctx->ac.builder, alpha, ctx->ac.i32, "");
2067 }
2068
2069 static unsigned
2070 get_num_channels_from_data_format(unsigned data_format)
2071 {
2072 switch (data_format) {
2073 case V_008F0C_BUF_DATA_FORMAT_8:
2074 case V_008F0C_BUF_DATA_FORMAT_16:
2075 case V_008F0C_BUF_DATA_FORMAT_32:
2076 return 1;
2077 case V_008F0C_BUF_DATA_FORMAT_8_8:
2078 case V_008F0C_BUF_DATA_FORMAT_16_16:
2079 case V_008F0C_BUF_DATA_FORMAT_32_32:
2080 return 2;
2081 case V_008F0C_BUF_DATA_FORMAT_10_11_11:
2082 case V_008F0C_BUF_DATA_FORMAT_11_11_10:
2083 case V_008F0C_BUF_DATA_FORMAT_32_32_32:
2084 return 3;
2085 case V_008F0C_BUF_DATA_FORMAT_8_8_8_8:
2086 case V_008F0C_BUF_DATA_FORMAT_10_10_10_2:
2087 case V_008F0C_BUF_DATA_FORMAT_2_10_10_10:
2088 case V_008F0C_BUF_DATA_FORMAT_16_16_16_16:
2089 case V_008F0C_BUF_DATA_FORMAT_32_32_32_32:
2090 return 4;
2091 default:
2092 break;
2093 }
2094
2095 return 4;
2096 }
2097
2098 static LLVMValueRef
2099 radv_fixup_vertex_input_fetches(struct radv_shader_context *ctx,
2100 LLVMValueRef value,
2101 unsigned num_channels,
2102 bool is_float)
2103 {
2104 LLVMValueRef zero = is_float ? ctx->ac.f32_0 : ctx->ac.i32_0;
2105 LLVMValueRef one = is_float ? ctx->ac.f32_1 : ctx->ac.i32_1;
2106 LLVMValueRef chan[4];
2107
2108 if (LLVMGetTypeKind(LLVMTypeOf(value)) == LLVMVectorTypeKind) {
2109 unsigned vec_size = LLVMGetVectorSize(LLVMTypeOf(value));
2110
2111 if (num_channels == 4 && num_channels == vec_size)
2112 return value;
2113
2114 num_channels = MIN2(num_channels, vec_size);
2115
2116 for (unsigned i = 0; i < num_channels; i++)
2117 chan[i] = ac_llvm_extract_elem(&ctx->ac, value, i);
2118 } else {
2119 if (num_channels) {
2120 assert(num_channels == 1);
2121 chan[0] = value;
2122 }
2123 }
2124
2125 for (unsigned i = num_channels; i < 4; i++) {
2126 chan[i] = i == 3 ? one : zero;
2127 chan[i] = ac_to_integer(&ctx->ac, chan[i]);
2128 }
2129
2130 return ac_build_gather_values(&ctx->ac, chan, 4);
2131 }
2132
2133 static void
2134 handle_vs_input_decl(struct radv_shader_context *ctx,
2135 struct nir_variable *variable)
2136 {
2137 LLVMValueRef t_list_ptr = ctx->vertex_buffers;
2138 LLVMValueRef t_offset;
2139 LLVMValueRef t_list;
2140 LLVMValueRef input;
2141 LLVMValueRef buffer_index;
2142 unsigned attrib_count = glsl_count_attribute_slots(variable->type, true);
2143 uint8_t input_usage_mask =
2144 ctx->shader_info->vs.input_usage_mask[variable->data.location];
2145 unsigned num_input_channels = util_last_bit(input_usage_mask);
2146
2147 variable->data.driver_location = variable->data.location * 4;
2148
2149 enum glsl_base_type type = glsl_get_base_type(variable->type);
2150 for (unsigned i = 0; i < attrib_count; ++i) {
2151 LLVMValueRef output[4];
2152 unsigned attrib_index = variable->data.location + i - VERT_ATTRIB_GENERIC0;
2153 unsigned attrib_format = ctx->options->key.vs.vertex_attribute_formats[attrib_index];
2154 unsigned data_format = attrib_format & 0x0f;
2155 unsigned num_format = (attrib_format >> 4) & 0x07;
2156 bool is_float = num_format != V_008F0C_BUF_NUM_FORMAT_UINT &&
2157 num_format != V_008F0C_BUF_NUM_FORMAT_SINT;
2158
2159 if (ctx->options->key.vs.instance_rate_inputs & (1u << attrib_index)) {
2160 uint32_t divisor = ctx->options->key.vs.instance_rate_divisors[attrib_index];
2161
2162 if (divisor) {
2163 buffer_index = ctx->abi.instance_id;
2164
2165 if (divisor != 1) {
2166 buffer_index = LLVMBuildUDiv(ctx->ac.builder, buffer_index,
2167 LLVMConstInt(ctx->ac.i32, divisor, 0), "");
2168 }
2169 } else {
2170 buffer_index = ctx->ac.i32_0;
2171 }
2172
2173 buffer_index = LLVMBuildAdd(ctx->ac.builder, ctx->abi.start_instance, buffer_index, "");
2174 } else
2175 buffer_index = LLVMBuildAdd(ctx->ac.builder, ctx->abi.vertex_id,
2176 ctx->abi.base_vertex, "");
2177
2178 /* Adjust the number of channels to load based on the vertex
2179 * attribute format.
2180 */
2181 unsigned num_format_channels = get_num_channels_from_data_format(data_format);
2182 unsigned num_channels = MIN2(num_input_channels, num_format_channels);
2183 unsigned attrib_binding = ctx->options->key.vs.vertex_attribute_bindings[attrib_index];
2184 unsigned attrib_offset = ctx->options->key.vs.vertex_attribute_offsets[attrib_index];
2185 unsigned attrib_stride = ctx->options->key.vs.vertex_attribute_strides[attrib_index];
2186
2187 if (ctx->options->key.vs.post_shuffle & (1 << attrib_index)) {
2188 /* Always load, at least, 3 channels for formats that
2189 * need to be shuffled because X<->Z.
2190 */
2191 num_channels = MAX2(num_channels, 3);
2192 }
2193
2194 if (attrib_stride != 0 && attrib_offset > attrib_stride) {
2195 LLVMValueRef buffer_offset =
2196 LLVMConstInt(ctx->ac.i32,
2197 attrib_offset / attrib_stride, false);
2198
2199 buffer_index = LLVMBuildAdd(ctx->ac.builder,
2200 buffer_index,
2201 buffer_offset, "");
2202
2203 attrib_offset = attrib_offset % attrib_stride;
2204 }
2205
2206 t_offset = LLVMConstInt(ctx->ac.i32, attrib_binding, false);
2207 t_list = ac_build_load_to_sgpr(&ctx->ac, t_list_ptr, t_offset);
2208
2209 input = ac_build_struct_tbuffer_load(&ctx->ac, t_list,
2210 buffer_index,
2211 LLVMConstInt(ctx->ac.i32, attrib_offset, false),
2212 ctx->ac.i32_0, ctx->ac.i32_0,
2213 num_channels,
2214 data_format, num_format, 0, true);
2215
2216 if (ctx->options->key.vs.post_shuffle & (1 << attrib_index)) {
2217 LLVMValueRef c[4];
2218 c[0] = ac_llvm_extract_elem(&ctx->ac, input, 2);
2219 c[1] = ac_llvm_extract_elem(&ctx->ac, input, 1);
2220 c[2] = ac_llvm_extract_elem(&ctx->ac, input, 0);
2221 c[3] = ac_llvm_extract_elem(&ctx->ac, input, 3);
2222
2223 input = ac_build_gather_values(&ctx->ac, c, 4);
2224 }
2225
2226 input = radv_fixup_vertex_input_fetches(ctx, input, num_channels,
2227 is_float);
2228
2229 for (unsigned chan = 0; chan < 4; chan++) {
2230 LLVMValueRef llvm_chan = LLVMConstInt(ctx->ac.i32, chan, false);
2231 output[chan] = LLVMBuildExtractElement(ctx->ac.builder, input, llvm_chan, "");
2232 if (type == GLSL_TYPE_FLOAT16) {
2233 output[chan] = LLVMBuildBitCast(ctx->ac.builder, output[chan], ctx->ac.f32, "");
2234 output[chan] = LLVMBuildFPTrunc(ctx->ac.builder, output[chan], ctx->ac.f16, "");
2235 }
2236 }
2237
2238 unsigned alpha_adjust = (ctx->options->key.vs.alpha_adjust >> (attrib_index * 2)) & 3;
2239 output[3] = adjust_vertex_fetch_alpha(ctx, alpha_adjust, output[3]);
2240
2241 for (unsigned chan = 0; chan < 4; chan++) {
2242 output[chan] = ac_to_integer(&ctx->ac, output[chan]);
2243 if (type == GLSL_TYPE_UINT16 || type == GLSL_TYPE_INT16)
2244 output[chan] = LLVMBuildTrunc(ctx->ac.builder, output[chan], ctx->ac.i16, "");
2245
2246 ctx->inputs[ac_llvm_reg_index_soa(variable->data.location + i, chan)] = output[chan];
2247 }
2248 }
2249 }
2250
2251 static void
2252 handle_vs_inputs(struct radv_shader_context *ctx,
2253 struct nir_shader *nir) {
2254 nir_foreach_variable(variable, &nir->inputs)
2255 handle_vs_input_decl(ctx, variable);
2256 }
2257
2258 static void
2259 prepare_interp_optimize(struct radv_shader_context *ctx,
2260 struct nir_shader *nir)
2261 {
2262 bool uses_center = false;
2263 bool uses_centroid = false;
2264 nir_foreach_variable(variable, &nir->inputs) {
2265 if (glsl_get_base_type(glsl_without_array(variable->type)) != GLSL_TYPE_FLOAT ||
2266 variable->data.sample)
2267 continue;
2268
2269 if (variable->data.centroid)
2270 uses_centroid = true;
2271 else
2272 uses_center = true;
2273 }
2274
2275 if (uses_center && uses_centroid) {
2276 LLVMValueRef sel = LLVMBuildICmp(ctx->ac.builder, LLVMIntSLT, ctx->abi.prim_mask, ctx->ac.i32_0, "");
2277 ctx->abi.persp_centroid = LLVMBuildSelect(ctx->ac.builder, sel, ctx->abi.persp_center, ctx->abi.persp_centroid, "");
2278 ctx->abi.linear_centroid = LLVMBuildSelect(ctx->ac.builder, sel, ctx->abi.linear_center, ctx->abi.linear_centroid, "");
2279 }
2280 }
2281
2282 static void
2283 scan_shader_output_decl(struct radv_shader_context *ctx,
2284 struct nir_variable *variable,
2285 struct nir_shader *shader,
2286 gl_shader_stage stage)
2287 {
2288 int idx = variable->data.location + variable->data.index;
2289 unsigned attrib_count = glsl_count_attribute_slots(variable->type, false);
2290 uint64_t mask_attribs;
2291
2292 variable->data.driver_location = idx * 4;
2293
2294 /* tess ctrl has it's own load/store paths for outputs */
2295 if (stage == MESA_SHADER_TESS_CTRL)
2296 return;
2297
2298 if (variable->data.compact) {
2299 unsigned component_count = variable->data.location_frac +
2300 glsl_get_length(variable->type);
2301 attrib_count = (component_count + 3) / 4;
2302 }
2303
2304 mask_attribs = ((1ull << attrib_count) - 1) << idx;
2305
2306 ctx->output_mask |= mask_attribs;
2307 }
2308
2309
2310 /* Initialize arguments for the shader export intrinsic */
2311 static void
2312 si_llvm_init_export_args(struct radv_shader_context *ctx,
2313 LLVMValueRef *values,
2314 unsigned enabled_channels,
2315 unsigned target,
2316 struct ac_export_args *args)
2317 {
2318 /* Specify the channels that are enabled. */
2319 args->enabled_channels = enabled_channels;
2320
2321 /* Specify whether the EXEC mask represents the valid mask */
2322 args->valid_mask = 0;
2323
2324 /* Specify whether this is the last export */
2325 args->done = 0;
2326
2327 /* Specify the target we are exporting */
2328 args->target = target;
2329
2330 args->compr = false;
2331 args->out[0] = LLVMGetUndef(ctx->ac.f32);
2332 args->out[1] = LLVMGetUndef(ctx->ac.f32);
2333 args->out[2] = LLVMGetUndef(ctx->ac.f32);
2334 args->out[3] = LLVMGetUndef(ctx->ac.f32);
2335
2336 if (!values)
2337 return;
2338
2339 bool is_16bit = ac_get_type_size(LLVMTypeOf(values[0])) == 2;
2340 if (ctx->stage == MESA_SHADER_FRAGMENT) {
2341 unsigned index = target - V_008DFC_SQ_EXP_MRT;
2342 unsigned col_format = (ctx->options->key.fs.col_format >> (4 * index)) & 0xf;
2343 bool is_int8 = (ctx->options->key.fs.is_int8 >> index) & 1;
2344 bool is_int10 = (ctx->options->key.fs.is_int10 >> index) & 1;
2345 unsigned chan;
2346
2347 LLVMValueRef (*packf)(struct ac_llvm_context *ctx, LLVMValueRef args[2]) = NULL;
2348 LLVMValueRef (*packi)(struct ac_llvm_context *ctx, LLVMValueRef args[2],
2349 unsigned bits, bool hi) = NULL;
2350
2351 switch(col_format) {
2352 case V_028714_SPI_SHADER_ZERO:
2353 args->enabled_channels = 0; /* writemask */
2354 args->target = V_008DFC_SQ_EXP_NULL;
2355 break;
2356
2357 case V_028714_SPI_SHADER_32_R:
2358 args->enabled_channels = 1;
2359 args->out[0] = values[0];
2360 break;
2361
2362 case V_028714_SPI_SHADER_32_GR:
2363 args->enabled_channels = 0x3;
2364 args->out[0] = values[0];
2365 args->out[1] = values[1];
2366 break;
2367
2368 case V_028714_SPI_SHADER_32_AR:
2369 if (ctx->ac.chip_class >= GFX10) {
2370 args->enabled_channels = 0x3;
2371 args->out[0] = values[0];
2372 args->out[1] = values[3];
2373 } else {
2374 args->enabled_channels = 0x9;
2375 args->out[0] = values[0];
2376 args->out[3] = values[3];
2377 }
2378 break;
2379
2380 case V_028714_SPI_SHADER_FP16_ABGR:
2381 args->enabled_channels = 0x5;
2382 packf = ac_build_cvt_pkrtz_f16;
2383 if (is_16bit) {
2384 for (unsigned chan = 0; chan < 4; chan++)
2385 values[chan] = LLVMBuildFPExt(ctx->ac.builder,
2386 values[chan],
2387 ctx->ac.f32, "");
2388 }
2389 break;
2390
2391 case V_028714_SPI_SHADER_UNORM16_ABGR:
2392 args->enabled_channels = 0x5;
2393 packf = ac_build_cvt_pknorm_u16;
2394 break;
2395
2396 case V_028714_SPI_SHADER_SNORM16_ABGR:
2397 args->enabled_channels = 0x5;
2398 packf = ac_build_cvt_pknorm_i16;
2399 break;
2400
2401 case V_028714_SPI_SHADER_UINT16_ABGR:
2402 args->enabled_channels = 0x5;
2403 packi = ac_build_cvt_pk_u16;
2404 if (is_16bit) {
2405 for (unsigned chan = 0; chan < 4; chan++)
2406 values[chan] = LLVMBuildZExt(ctx->ac.builder,
2407 ac_to_integer(&ctx->ac, values[chan]),
2408 ctx->ac.i32, "");
2409 }
2410 break;
2411
2412 case V_028714_SPI_SHADER_SINT16_ABGR:
2413 args->enabled_channels = 0x5;
2414 packi = ac_build_cvt_pk_i16;
2415 if (is_16bit) {
2416 for (unsigned chan = 0; chan < 4; chan++)
2417 values[chan] = LLVMBuildSExt(ctx->ac.builder,
2418 ac_to_integer(&ctx->ac, values[chan]),
2419 ctx->ac.i32, "");
2420 }
2421 break;
2422
2423 default:
2424 case V_028714_SPI_SHADER_32_ABGR:
2425 memcpy(&args->out[0], values, sizeof(values[0]) * 4);
2426 break;
2427 }
2428
2429 /* Pack f16 or norm_i16/u16. */
2430 if (packf) {
2431 for (chan = 0; chan < 2; chan++) {
2432 LLVMValueRef pack_args[2] = {
2433 values[2 * chan],
2434 values[2 * chan + 1]
2435 };
2436 LLVMValueRef packed;
2437
2438 packed = packf(&ctx->ac, pack_args);
2439 args->out[chan] = ac_to_float(&ctx->ac, packed);
2440 }
2441 args->compr = 1; /* COMPR flag */
2442 }
2443
2444 /* Pack i16/u16. */
2445 if (packi) {
2446 for (chan = 0; chan < 2; chan++) {
2447 LLVMValueRef pack_args[2] = {
2448 ac_to_integer(&ctx->ac, values[2 * chan]),
2449 ac_to_integer(&ctx->ac, values[2 * chan + 1])
2450 };
2451 LLVMValueRef packed;
2452
2453 packed = packi(&ctx->ac, pack_args,
2454 is_int8 ? 8 : is_int10 ? 10 : 16,
2455 chan == 1);
2456 args->out[chan] = ac_to_float(&ctx->ac, packed);
2457 }
2458 args->compr = 1; /* COMPR flag */
2459 }
2460 return;
2461 }
2462
2463 if (is_16bit) {
2464 for (unsigned chan = 0; chan < 4; chan++) {
2465 values[chan] = LLVMBuildBitCast(ctx->ac.builder, values[chan], ctx->ac.i16, "");
2466 args->out[chan] = LLVMBuildZExt(ctx->ac.builder, values[chan], ctx->ac.i32, "");
2467 }
2468 } else
2469 memcpy(&args->out[0], values, sizeof(values[0]) * 4);
2470
2471 for (unsigned i = 0; i < 4; ++i)
2472 args->out[i] = ac_to_float(&ctx->ac, args->out[i]);
2473 }
2474
2475 static void
2476 radv_export_param(struct radv_shader_context *ctx, unsigned index,
2477 LLVMValueRef *values, unsigned enabled_channels)
2478 {
2479 struct ac_export_args args;
2480
2481 si_llvm_init_export_args(ctx, values, enabled_channels,
2482 V_008DFC_SQ_EXP_PARAM + index, &args);
2483 ac_build_export(&ctx->ac, &args);
2484 }
2485
2486 static LLVMValueRef
2487 radv_load_output(struct radv_shader_context *ctx, unsigned index, unsigned chan)
2488 {
2489 LLVMValueRef output = ctx->abi.outputs[ac_llvm_reg_index_soa(index, chan)];
2490 return LLVMBuildLoad(ctx->ac.builder, output, "");
2491 }
2492
2493 static void
2494 radv_emit_stream_output(struct radv_shader_context *ctx,
2495 LLVMValueRef const *so_buffers,
2496 LLVMValueRef const *so_write_offsets,
2497 const struct radv_stream_output *output,
2498 struct radv_shader_output_values *shader_out)
2499 {
2500 unsigned num_comps = util_bitcount(output->component_mask);
2501 unsigned buf = output->buffer;
2502 unsigned offset = output->offset;
2503 unsigned start;
2504 LLVMValueRef out[4];
2505
2506 assert(num_comps && num_comps <= 4);
2507 if (!num_comps || num_comps > 4)
2508 return;
2509
2510 /* Get the first component. */
2511 start = ffs(output->component_mask) - 1;
2512
2513 /* Load the output as int. */
2514 for (int i = 0; i < num_comps; i++) {
2515 out[i] = ac_to_integer(&ctx->ac, shader_out->values[start + i]);
2516 }
2517
2518 /* Pack the output. */
2519 LLVMValueRef vdata = NULL;
2520
2521 switch (num_comps) {
2522 case 1: /* as i32 */
2523 vdata = out[0];
2524 break;
2525 case 2: /* as v2i32 */
2526 case 3: /* as v4i32 (aligned to 4) */
2527 out[3] = LLVMGetUndef(ctx->ac.i32);
2528 /* fall through */
2529 case 4: /* as v4i32 */
2530 vdata = ac_build_gather_values(&ctx->ac, out,
2531 !ac_has_vec3_support(ctx->ac.chip_class, false) ?
2532 util_next_power_of_two(num_comps) :
2533 num_comps);
2534 break;
2535 }
2536
2537 ac_build_buffer_store_dword(&ctx->ac, so_buffers[buf],
2538 vdata, num_comps, so_write_offsets[buf],
2539 ctx->ac.i32_0, offset,
2540 ac_glc | ac_slc, false);
2541 }
2542
2543 static void
2544 radv_emit_streamout(struct radv_shader_context *ctx, unsigned stream)
2545 {
2546 int i;
2547
2548 /* Get bits [22:16], i.e. (so_param >> 16) & 127; */
2549 assert(ctx->streamout_config);
2550 LLVMValueRef so_vtx_count =
2551 ac_build_bfe(&ctx->ac, ctx->streamout_config,
2552 LLVMConstInt(ctx->ac.i32, 16, false),
2553 LLVMConstInt(ctx->ac.i32, 7, false), false);
2554
2555 LLVMValueRef tid = ac_get_thread_id(&ctx->ac);
2556
2557 /* can_emit = tid < so_vtx_count; */
2558 LLVMValueRef can_emit = LLVMBuildICmp(ctx->ac.builder, LLVMIntULT,
2559 tid, so_vtx_count, "");
2560
2561 /* Emit the streamout code conditionally. This actually avoids
2562 * out-of-bounds buffer access. The hw tells us via the SGPR
2563 * (so_vtx_count) which threads are allowed to emit streamout data.
2564 */
2565 ac_build_ifcc(&ctx->ac, can_emit, 6501);
2566 {
2567 /* The buffer offset is computed as follows:
2568 * ByteOffset = streamout_offset[buffer_id]*4 +
2569 * (streamout_write_index + thread_id)*stride[buffer_id] +
2570 * attrib_offset
2571 */
2572 LLVMValueRef so_write_index = ctx->streamout_write_idx;
2573
2574 /* Compute (streamout_write_index + thread_id). */
2575 so_write_index =
2576 LLVMBuildAdd(ctx->ac.builder, so_write_index, tid, "");
2577
2578 /* Load the descriptor and compute the write offset for each
2579 * enabled buffer.
2580 */
2581 LLVMValueRef so_write_offset[4] = {};
2582 LLVMValueRef so_buffers[4] = {};
2583 LLVMValueRef buf_ptr = ctx->streamout_buffers;
2584
2585 for (i = 0; i < 4; i++) {
2586 uint16_t stride = ctx->shader_info->so.strides[i];
2587
2588 if (!stride)
2589 continue;
2590
2591 LLVMValueRef offset =
2592 LLVMConstInt(ctx->ac.i32, i, false);
2593
2594 so_buffers[i] = ac_build_load_to_sgpr(&ctx->ac,
2595 buf_ptr, offset);
2596
2597 LLVMValueRef so_offset = ctx->streamout_offset[i];
2598
2599 so_offset = LLVMBuildMul(ctx->ac.builder, so_offset,
2600 LLVMConstInt(ctx->ac.i32, 4, false), "");
2601
2602 so_write_offset[i] =
2603 ac_build_imad(&ctx->ac, so_write_index,
2604 LLVMConstInt(ctx->ac.i32,
2605 stride * 4, false),
2606 so_offset);
2607 }
2608
2609 /* Write streamout data. */
2610 for (i = 0; i < ctx->shader_info->so.num_outputs; i++) {
2611 struct radv_shader_output_values shader_out = {};
2612 struct radv_stream_output *output =
2613 &ctx->shader_info->so.outputs[i];
2614
2615 if (stream != output->stream)
2616 continue;
2617
2618 for (int j = 0; j < 4; j++) {
2619 shader_out.values[j] =
2620 radv_load_output(ctx, output->location, j);
2621 }
2622
2623 radv_emit_stream_output(ctx, so_buffers,so_write_offset,
2624 output, &shader_out);
2625 }
2626 }
2627 ac_build_endif(&ctx->ac, 6501);
2628 }
2629
2630 static void
2631 radv_build_param_exports(struct radv_shader_context *ctx,
2632 struct radv_shader_output_values *outputs,
2633 unsigned noutput,
2634 struct radv_vs_output_info *outinfo,
2635 bool export_clip_dists)
2636 {
2637 unsigned param_count = 0;
2638
2639 for (unsigned i = 0; i < noutput; i++) {
2640 unsigned slot_name = outputs[i].slot_name;
2641 unsigned usage_mask = outputs[i].usage_mask;
2642
2643 if (slot_name != VARYING_SLOT_LAYER &&
2644 slot_name != VARYING_SLOT_PRIMITIVE_ID &&
2645 slot_name != VARYING_SLOT_CLIP_DIST0 &&
2646 slot_name != VARYING_SLOT_CLIP_DIST1 &&
2647 slot_name < VARYING_SLOT_VAR0)
2648 continue;
2649
2650 if ((slot_name == VARYING_SLOT_CLIP_DIST0 ||
2651 slot_name == VARYING_SLOT_CLIP_DIST1) && !export_clip_dists)
2652 continue;
2653
2654 radv_export_param(ctx, param_count, outputs[i].values, usage_mask);
2655
2656 assert(i < ARRAY_SIZE(outinfo->vs_output_param_offset));
2657 outinfo->vs_output_param_offset[slot_name] = param_count++;
2658 }
2659
2660 outinfo->param_exports = param_count;
2661 }
2662
2663 /* Generate export instructions for hardware VS shader stage or NGG GS stage
2664 * (position and parameter data only).
2665 */
2666 static void
2667 radv_llvm_export_vs(struct radv_shader_context *ctx,
2668 struct radv_shader_output_values *outputs,
2669 unsigned noutput,
2670 struct radv_vs_output_info *outinfo,
2671 bool export_clip_dists)
2672 {
2673 LLVMValueRef psize_value = NULL, layer_value = NULL, viewport_value = NULL;
2674 struct ac_export_args pos_args[4] = {};
2675 unsigned pos_idx, index;
2676 int i;
2677
2678 /* Build position exports */
2679 for (i = 0; i < noutput; i++) {
2680 switch (outputs[i].slot_name) {
2681 case VARYING_SLOT_POS:
2682 si_llvm_init_export_args(ctx, outputs[i].values, 0xf,
2683 V_008DFC_SQ_EXP_POS, &pos_args[0]);
2684 break;
2685 case VARYING_SLOT_PSIZ:
2686 psize_value = outputs[i].values[0];
2687 break;
2688 case VARYING_SLOT_LAYER:
2689 layer_value = outputs[i].values[0];
2690 break;
2691 case VARYING_SLOT_VIEWPORT:
2692 viewport_value = outputs[i].values[0];
2693 break;
2694 case VARYING_SLOT_CLIP_DIST0:
2695 case VARYING_SLOT_CLIP_DIST1:
2696 index = 2 + outputs[i].slot_index;
2697 si_llvm_init_export_args(ctx, outputs[i].values, 0xf,
2698 V_008DFC_SQ_EXP_POS + index,
2699 &pos_args[index]);
2700 break;
2701 default:
2702 break;
2703 }
2704 }
2705
2706 /* We need to add the position output manually if it's missing. */
2707 if (!pos_args[0].out[0]) {
2708 pos_args[0].enabled_channels = 0xf; /* writemask */
2709 pos_args[0].valid_mask = 0; /* EXEC mask */
2710 pos_args[0].done = 0; /* last export? */
2711 pos_args[0].target = V_008DFC_SQ_EXP_POS;
2712 pos_args[0].compr = 0; /* COMPR flag */
2713 pos_args[0].out[0] = ctx->ac.f32_0; /* X */
2714 pos_args[0].out[1] = ctx->ac.f32_0; /* Y */
2715 pos_args[0].out[2] = ctx->ac.f32_0; /* Z */
2716 pos_args[0].out[3] = ctx->ac.f32_1; /* W */
2717 }
2718
2719 if (outinfo->writes_pointsize ||
2720 outinfo->writes_layer ||
2721 outinfo->writes_viewport_index) {
2722 pos_args[1].enabled_channels = ((outinfo->writes_pointsize == true ? 1 : 0) |
2723 (outinfo->writes_layer == true ? 4 : 0));
2724 pos_args[1].valid_mask = 0;
2725 pos_args[1].done = 0;
2726 pos_args[1].target = V_008DFC_SQ_EXP_POS + 1;
2727 pos_args[1].compr = 0;
2728 pos_args[1].out[0] = ctx->ac.f32_0; /* X */
2729 pos_args[1].out[1] = ctx->ac.f32_0; /* Y */
2730 pos_args[1].out[2] = ctx->ac.f32_0; /* Z */
2731 pos_args[1].out[3] = ctx->ac.f32_0; /* W */
2732
2733 if (outinfo->writes_pointsize == true)
2734 pos_args[1].out[0] = psize_value;
2735 if (outinfo->writes_layer == true)
2736 pos_args[1].out[2] = layer_value;
2737 if (outinfo->writes_viewport_index == true) {
2738 if (ctx->options->chip_class >= GFX9) {
2739 /* GFX9 has the layer in out.z[10:0] and the viewport
2740 * index in out.z[19:16].
2741 */
2742 LLVMValueRef v = viewport_value;
2743 v = ac_to_integer(&ctx->ac, v);
2744 v = LLVMBuildShl(ctx->ac.builder, v,
2745 LLVMConstInt(ctx->ac.i32, 16, false),
2746 "");
2747 v = LLVMBuildOr(ctx->ac.builder, v,
2748 ac_to_integer(&ctx->ac, pos_args[1].out[2]), "");
2749
2750 pos_args[1].out[2] = ac_to_float(&ctx->ac, v);
2751 pos_args[1].enabled_channels |= 1 << 2;
2752 } else {
2753 pos_args[1].out[3] = viewport_value;
2754 pos_args[1].enabled_channels |= 1 << 3;
2755 }
2756 }
2757 }
2758
2759 for (i = 0; i < 4; i++) {
2760 if (pos_args[i].out[0])
2761 outinfo->pos_exports++;
2762 }
2763
2764 /* Navi10-14 skip POS0 exports if EXEC=0 and DONE=0, causing a hang.
2765 * Setting valid_mask=1 prevents it and has no other effect.
2766 */
2767 if (ctx->ac.family == CHIP_NAVI10 ||
2768 ctx->ac.family == CHIP_NAVI12 ||
2769 ctx->ac.family == CHIP_NAVI14)
2770 pos_args[0].valid_mask = 1;
2771
2772 pos_idx = 0;
2773 for (i = 0; i < 4; i++) {
2774 if (!pos_args[i].out[0])
2775 continue;
2776
2777 /* Specify the target we are exporting */
2778 pos_args[i].target = V_008DFC_SQ_EXP_POS + pos_idx++;
2779
2780 if (pos_idx == outinfo->pos_exports)
2781 /* Specify that this is the last export */
2782 pos_args[i].done = 1;
2783
2784 ac_build_export(&ctx->ac, &pos_args[i]);
2785 }
2786
2787 /* Build parameter exports */
2788 radv_build_param_exports(ctx, outputs, noutput, outinfo, export_clip_dists);
2789 }
2790
2791 static void
2792 handle_vs_outputs_post(struct radv_shader_context *ctx,
2793 bool export_prim_id,
2794 bool export_clip_dists,
2795 struct radv_vs_output_info *outinfo)
2796 {
2797 struct radv_shader_output_values *outputs;
2798 unsigned noutput = 0;
2799
2800 if (ctx->options->key.has_multiview_view_index) {
2801 LLVMValueRef* tmp_out = &ctx->abi.outputs[ac_llvm_reg_index_soa(VARYING_SLOT_LAYER, 0)];
2802 if(!*tmp_out) {
2803 for(unsigned i = 0; i < 4; ++i)
2804 ctx->abi.outputs[ac_llvm_reg_index_soa(VARYING_SLOT_LAYER, i)] =
2805 ac_build_alloca_undef(&ctx->ac, ctx->ac.f32, "");
2806 }
2807
2808 LLVMBuildStore(ctx->ac.builder, ac_to_float(&ctx->ac, ctx->abi.view_index), *tmp_out);
2809 ctx->output_mask |= 1ull << VARYING_SLOT_LAYER;
2810 }
2811
2812 memset(outinfo->vs_output_param_offset, AC_EXP_PARAM_UNDEFINED,
2813 sizeof(outinfo->vs_output_param_offset));
2814 outinfo->pos_exports = 0;
2815
2816 if (!ctx->options->use_ngg_streamout &&
2817 ctx->shader_info->so.num_outputs &&
2818 !ctx->is_gs_copy_shader) {
2819 /* The GS copy shader emission already emits streamout. */
2820 radv_emit_streamout(ctx, 0);
2821 }
2822
2823 /* Allocate a temporary array for the output values. */
2824 unsigned num_outputs = util_bitcount64(ctx->output_mask) + export_prim_id;
2825 outputs = malloc(num_outputs * sizeof(outputs[0]));
2826
2827 for (unsigned i = 0; i < AC_LLVM_MAX_OUTPUTS; ++i) {
2828 if (!(ctx->output_mask & (1ull << i)))
2829 continue;
2830
2831 outputs[noutput].slot_name = i;
2832 outputs[noutput].slot_index = i == VARYING_SLOT_CLIP_DIST1;
2833
2834 if (ctx->stage == MESA_SHADER_VERTEX &&
2835 !ctx->is_gs_copy_shader) {
2836 outputs[noutput].usage_mask =
2837 ctx->shader_info->vs.output_usage_mask[i];
2838 } else if (ctx->stage == MESA_SHADER_TESS_EVAL) {
2839 outputs[noutput].usage_mask =
2840 ctx->shader_info->tes.output_usage_mask[i];
2841 } else {
2842 assert(ctx->is_gs_copy_shader);
2843 outputs[noutput].usage_mask =
2844 ctx->shader_info->gs.output_usage_mask[i];
2845 }
2846
2847 for (unsigned j = 0; j < 4; j++) {
2848 outputs[noutput].values[j] =
2849 ac_to_float(&ctx->ac, radv_load_output(ctx, i, j));
2850 }
2851
2852 noutput++;
2853 }
2854
2855 /* Export PrimitiveID. */
2856 if (export_prim_id) {
2857 outputs[noutput].slot_name = VARYING_SLOT_PRIMITIVE_ID;
2858 outputs[noutput].slot_index = 0;
2859 outputs[noutput].usage_mask = 0x1;
2860 outputs[noutput].values[0] = ctx->vs_prim_id;
2861 for (unsigned j = 1; j < 4; j++)
2862 outputs[noutput].values[j] = ctx->ac.f32_0;
2863 noutput++;
2864 }
2865
2866 radv_llvm_export_vs(ctx, outputs, noutput, outinfo, export_clip_dists);
2867
2868 free(outputs);
2869 }
2870
2871 static void
2872 handle_es_outputs_post(struct radv_shader_context *ctx,
2873 struct radv_es_output_info *outinfo)
2874 {
2875 int j;
2876 LLVMValueRef lds_base = NULL;
2877
2878 if (ctx->ac.chip_class >= GFX9) {
2879 unsigned itemsize_dw = outinfo->esgs_itemsize / 4;
2880 LLVMValueRef vertex_idx = ac_get_thread_id(&ctx->ac);
2881 LLVMValueRef wave_idx = ac_unpack_param(&ctx->ac, ctx->merged_wave_info, 24, 4);
2882 vertex_idx = LLVMBuildOr(ctx->ac.builder, vertex_idx,
2883 LLVMBuildMul(ctx->ac.builder, wave_idx,
2884 LLVMConstInt(ctx->ac.i32,
2885 ctx->ac.wave_size, false), ""), "");
2886 lds_base = LLVMBuildMul(ctx->ac.builder, vertex_idx,
2887 LLVMConstInt(ctx->ac.i32, itemsize_dw, 0), "");
2888 }
2889
2890 for (unsigned i = 0; i < AC_LLVM_MAX_OUTPUTS; ++i) {
2891 LLVMValueRef dw_addr = NULL;
2892 LLVMValueRef *out_ptr = &ctx->abi.outputs[i * 4];
2893 unsigned output_usage_mask;
2894 int param_index;
2895
2896 if (!(ctx->output_mask & (1ull << i)))
2897 continue;
2898
2899 if (ctx->stage == MESA_SHADER_VERTEX) {
2900 output_usage_mask =
2901 ctx->shader_info->vs.output_usage_mask[i];
2902 } else {
2903 assert(ctx->stage == MESA_SHADER_TESS_EVAL);
2904 output_usage_mask =
2905 ctx->shader_info->tes.output_usage_mask[i];
2906 }
2907
2908 param_index = shader_io_get_unique_index(i);
2909
2910 if (lds_base) {
2911 dw_addr = LLVMBuildAdd(ctx->ac.builder, lds_base,
2912 LLVMConstInt(ctx->ac.i32, param_index * 4, false),
2913 "");
2914 }
2915
2916 for (j = 0; j < 4; j++) {
2917 if (!(output_usage_mask & (1 << j)))
2918 continue;
2919
2920 LLVMValueRef out_val = LLVMBuildLoad(ctx->ac.builder, out_ptr[j], "");
2921 out_val = ac_to_integer(&ctx->ac, out_val);
2922 out_val = LLVMBuildZExtOrBitCast(ctx->ac.builder, out_val, ctx->ac.i32, "");
2923
2924 if (ctx->ac.chip_class >= GFX9) {
2925 LLVMValueRef dw_addr_offset =
2926 LLVMBuildAdd(ctx->ac.builder, dw_addr,
2927 LLVMConstInt(ctx->ac.i32,
2928 j, false), "");
2929
2930 ac_lds_store(&ctx->ac, dw_addr_offset, out_val);
2931 } else {
2932 ac_build_buffer_store_dword(&ctx->ac,
2933 ctx->esgs_ring,
2934 out_val, 1,
2935 NULL, ctx->es2gs_offset,
2936 (4 * param_index + j) * 4,
2937 ac_glc | ac_slc, true);
2938 }
2939 }
2940 }
2941 }
2942
2943 static void
2944 handle_ls_outputs_post(struct radv_shader_context *ctx)
2945 {
2946 LLVMValueRef vertex_id = ctx->rel_auto_id;
2947 uint32_t num_tcs_inputs = util_last_bit64(ctx->shader_info->vs.ls_outputs_written);
2948 LLVMValueRef vertex_dw_stride = LLVMConstInt(ctx->ac.i32, num_tcs_inputs * 4, false);
2949 LLVMValueRef base_dw_addr = LLVMBuildMul(ctx->ac.builder, vertex_id,
2950 vertex_dw_stride, "");
2951
2952 for (unsigned i = 0; i < AC_LLVM_MAX_OUTPUTS; ++i) {
2953 LLVMValueRef *out_ptr = &ctx->abi.outputs[i * 4];
2954
2955 if (!(ctx->output_mask & (1ull << i)))
2956 continue;
2957
2958 int param = shader_io_get_unique_index(i);
2959 LLVMValueRef dw_addr = LLVMBuildAdd(ctx->ac.builder, base_dw_addr,
2960 LLVMConstInt(ctx->ac.i32, param * 4, false),
2961 "");
2962 for (unsigned j = 0; j < 4; j++) {
2963 LLVMValueRef value = LLVMBuildLoad(ctx->ac.builder, out_ptr[j], "");
2964 value = ac_to_integer(&ctx->ac, value);
2965 value = LLVMBuildZExtOrBitCast(ctx->ac.builder, value, ctx->ac.i32, "");
2966 ac_lds_store(&ctx->ac, dw_addr, value);
2967 dw_addr = LLVMBuildAdd(ctx->ac.builder, dw_addr, ctx->ac.i32_1, "");
2968 }
2969 }
2970 }
2971
2972 static LLVMValueRef get_wave_id_in_tg(struct radv_shader_context *ctx)
2973 {
2974 return ac_unpack_param(&ctx->ac, ctx->merged_wave_info, 24, 4);
2975 }
2976
2977 static LLVMValueRef get_tgsize(struct radv_shader_context *ctx)
2978 {
2979 return ac_unpack_param(&ctx->ac, ctx->merged_wave_info, 28, 4);
2980 }
2981
2982 static LLVMValueRef get_thread_id_in_tg(struct radv_shader_context *ctx)
2983 {
2984 LLVMBuilderRef builder = ctx->ac.builder;
2985 LLVMValueRef tmp;
2986 tmp = LLVMBuildMul(builder, get_wave_id_in_tg(ctx),
2987 LLVMConstInt(ctx->ac.i32, ctx->ac.wave_size, false), "");
2988 return LLVMBuildAdd(builder, tmp, ac_get_thread_id(&ctx->ac), "");
2989 }
2990
2991 static LLVMValueRef ngg_get_vtx_cnt(struct radv_shader_context *ctx)
2992 {
2993 return ac_build_bfe(&ctx->ac, ctx->gs_tg_info,
2994 LLVMConstInt(ctx->ac.i32, 12, false),
2995 LLVMConstInt(ctx->ac.i32, 9, false),
2996 false);
2997 }
2998
2999 static LLVMValueRef ngg_get_prim_cnt(struct radv_shader_context *ctx)
3000 {
3001 return ac_build_bfe(&ctx->ac, ctx->gs_tg_info,
3002 LLVMConstInt(ctx->ac.i32, 22, false),
3003 LLVMConstInt(ctx->ac.i32, 9, false),
3004 false);
3005 }
3006
3007 static LLVMValueRef ngg_get_ordered_id(struct radv_shader_context *ctx)
3008 {
3009 return ac_build_bfe(&ctx->ac, ctx->gs_tg_info,
3010 ctx->ac.i32_0,
3011 LLVMConstInt(ctx->ac.i32, 11, false),
3012 false);
3013 }
3014
3015 static LLVMValueRef
3016 ngg_gs_get_vertex_storage(struct radv_shader_context *ctx)
3017 {
3018 unsigned num_outputs = util_bitcount64(ctx->output_mask);
3019
3020 if (ctx->options->key.has_multiview_view_index)
3021 num_outputs++;
3022
3023 LLVMTypeRef elements[2] = {
3024 LLVMArrayType(ctx->ac.i32, 4 * num_outputs),
3025 LLVMArrayType(ctx->ac.i8, 4),
3026 };
3027 LLVMTypeRef type = LLVMStructTypeInContext(ctx->ac.context, elements, 2, false);
3028 type = LLVMPointerType(LLVMArrayType(type, 0), AC_ADDR_SPACE_LDS);
3029 return LLVMBuildBitCast(ctx->ac.builder, ctx->gs_ngg_emit, type, "");
3030 }
3031
3032 /**
3033 * Return a pointer to the LDS storage reserved for the N'th vertex, where N
3034 * is in emit order; that is:
3035 * - during the epilogue, N is the threadidx (relative to the entire threadgroup)
3036 * - during vertex emit, i.e. while the API GS shader invocation is running,
3037 * N = threadidx * gs_max_out_vertices + emitidx
3038 *
3039 * Goals of the LDS memory layout:
3040 * 1. Eliminate bank conflicts on write for geometry shaders that have all emits
3041 * in uniform control flow
3042 * 2. Eliminate bank conflicts on read for export if, additionally, there is no
3043 * culling
3044 * 3. Agnostic to the number of waves (since we don't know it before compiling)
3045 * 4. Allow coalescing of LDS instructions (ds_write_b128 etc.)
3046 * 5. Avoid wasting memory.
3047 *
3048 * We use an AoS layout due to point 4 (this also helps point 3). In an AoS
3049 * layout, elimination of bank conflicts requires that each vertex occupy an
3050 * odd number of dwords. We use the additional dword to store the output stream
3051 * index as well as a flag to indicate whether this vertex ends a primitive
3052 * for rasterization.
3053 *
3054 * Swizzling is required to satisfy points 1 and 2 simultaneously.
3055 *
3056 * Vertices are stored in export order (gsthread * gs_max_out_vertices + emitidx).
3057 * Indices are swizzled in groups of 32, which ensures point 1 without
3058 * disturbing point 2.
3059 *
3060 * \return an LDS pointer to type {[N x i32], [4 x i8]}
3061 */
3062 static LLVMValueRef
3063 ngg_gs_vertex_ptr(struct radv_shader_context *ctx, LLVMValueRef vertexidx)
3064 {
3065 LLVMBuilderRef builder = ctx->ac.builder;
3066 LLVMValueRef storage = ngg_gs_get_vertex_storage(ctx);
3067
3068 /* gs_max_out_vertices = 2^(write_stride_2exp) * some odd number */
3069 unsigned write_stride_2exp = ffs(ctx->shader->info.gs.vertices_out) - 1;
3070 if (write_stride_2exp) {
3071 LLVMValueRef row =
3072 LLVMBuildLShr(builder, vertexidx,
3073 LLVMConstInt(ctx->ac.i32, 5, false), "");
3074 LLVMValueRef swizzle =
3075 LLVMBuildAnd(builder, row,
3076 LLVMConstInt(ctx->ac.i32, (1u << write_stride_2exp) - 1,
3077 false), "");
3078 vertexidx = LLVMBuildXor(builder, vertexidx, swizzle, "");
3079 }
3080
3081 return ac_build_gep0(&ctx->ac, storage, vertexidx);
3082 }
3083
3084 static LLVMValueRef
3085 ngg_gs_emit_vertex_ptr(struct radv_shader_context *ctx, LLVMValueRef gsthread,
3086 LLVMValueRef emitidx)
3087 {
3088 LLVMBuilderRef builder = ctx->ac.builder;
3089 LLVMValueRef tmp;
3090
3091 tmp = LLVMConstInt(ctx->ac.i32, ctx->shader->info.gs.vertices_out, false);
3092 tmp = LLVMBuildMul(builder, tmp, gsthread, "");
3093 const LLVMValueRef vertexidx = LLVMBuildAdd(builder, tmp, emitidx, "");
3094 return ngg_gs_vertex_ptr(ctx, vertexidx);
3095 }
3096
3097 /* Send GS Alloc Req message from the first wave of the group to SPI.
3098 * Message payload is:
3099 * - bits 0..10: vertices in group
3100 * - bits 12..22: primitives in group
3101 */
3102 static void build_sendmsg_gs_alloc_req(struct radv_shader_context *ctx,
3103 LLVMValueRef vtx_cnt,
3104 LLVMValueRef prim_cnt)
3105 {
3106 LLVMBuilderRef builder = ctx->ac.builder;
3107 LLVMValueRef tmp;
3108
3109 tmp = LLVMBuildICmp(builder, LLVMIntEQ, get_wave_id_in_tg(ctx), ctx->ac.i32_0, "");
3110 ac_build_ifcc(&ctx->ac, tmp, 5020);
3111
3112 tmp = LLVMBuildShl(builder, prim_cnt, LLVMConstInt(ctx->ac.i32, 12, false),"");
3113 tmp = LLVMBuildOr(builder, tmp, vtx_cnt, "");
3114 ac_build_sendmsg(&ctx->ac, AC_SENDMSG_GS_ALLOC_REQ, tmp);
3115
3116 ac_build_endif(&ctx->ac, 5020);
3117 }
3118
3119 struct ngg_prim {
3120 unsigned num_vertices;
3121 LLVMValueRef isnull;
3122 LLVMValueRef swap;
3123 LLVMValueRef index[3];
3124 LLVMValueRef edgeflag[3];
3125 };
3126
3127 static void build_export_prim(struct radv_shader_context *ctx,
3128 const struct ngg_prim *prim)
3129 {
3130 LLVMBuilderRef builder = ctx->ac.builder;
3131 struct ac_export_args args;
3132 LLVMValueRef vertices[3];
3133 LLVMValueRef odd, even;
3134 LLVMValueRef tmp;
3135
3136 tmp = LLVMBuildZExt(builder, prim->isnull, ctx->ac.i32, "");
3137 args.out[0] = LLVMBuildShl(builder, tmp, LLVMConstInt(ctx->ac.i32, 31, false), "");
3138
3139 for (unsigned i = 0; i < prim->num_vertices; ++i) {
3140 tmp = LLVMBuildZExt(builder, prim->edgeflag[i], ctx->ac.i32, "");
3141 tmp = LLVMBuildShl(builder, tmp,
3142 LLVMConstInt(ctx->ac.i32, 9, false), "");
3143 vertices[i] = LLVMBuildOr(builder, prim->index[i], tmp, "");
3144 }
3145
3146 switch (prim->num_vertices) {
3147 case 1:
3148 args.out[0] = LLVMBuildOr(builder, args.out[0], vertices[0], "");
3149 break;
3150 case 2:
3151 tmp = LLVMBuildShl(builder, vertices[1],
3152 LLVMConstInt(ctx->ac.i32, 10, false), "");
3153 tmp = LLVMBuildOr(builder, args.out[0], tmp, "");
3154 args.out[0] = LLVMBuildOr(builder, tmp, vertices[0], "");
3155 break;
3156 case 3:
3157 /* Swap vertices if needed to follow drawing order. */
3158 tmp = LLVMBuildShl(builder, vertices[2],
3159 LLVMConstInt(ctx->ac.i32, 20, false), "");
3160 even = LLVMBuildOr(builder, args.out[0], tmp, "");
3161 tmp = LLVMBuildShl(builder, vertices[1],
3162 LLVMConstInt(ctx->ac.i32, 10, false), "");
3163 even = LLVMBuildOr(builder, even, tmp, "");
3164 even = LLVMBuildOr(builder, even, vertices[0], "");
3165
3166 tmp = LLVMBuildShl(builder, vertices[1],
3167 LLVMConstInt(ctx->ac.i32, 20, false), "");
3168 odd = LLVMBuildOr(builder, args.out[0], tmp, "");
3169 tmp = LLVMBuildShl(builder, vertices[2],
3170 LLVMConstInt(ctx->ac.i32, 10, false), "");
3171 odd = LLVMBuildOr(builder, odd, tmp, "");
3172 odd = LLVMBuildOr(builder, odd, vertices[0], "");
3173
3174 args.out[0] = LLVMBuildSelect(builder, prim->swap, odd, even, "");
3175 break;
3176 default:
3177 unreachable("invalid number of vertices");
3178 }
3179
3180 args.out[0] = LLVMBuildBitCast(builder, args.out[0], ctx->ac.f32, "");
3181 args.out[1] = LLVMGetUndef(ctx->ac.f32);
3182 args.out[2] = LLVMGetUndef(ctx->ac.f32);
3183 args.out[3] = LLVMGetUndef(ctx->ac.f32);
3184
3185 args.target = V_008DFC_SQ_EXP_PRIM;
3186 args.enabled_channels = 1;
3187 args.done = true;
3188 args.valid_mask = false;
3189 args.compr = false;
3190
3191 ac_build_export(&ctx->ac, &args);
3192 }
3193
3194 static struct radv_stream_output *
3195 radv_get_stream_output_by_loc(struct radv_streamout_info *so, unsigned location)
3196 {
3197 for (unsigned i = 0; i < so->num_outputs; ++i) {
3198 if (so->outputs[i].location == location)
3199 return &so->outputs[i];
3200 }
3201
3202 return NULL;
3203 }
3204
3205 static void build_streamout_vertex(struct radv_shader_context *ctx,
3206 LLVMValueRef *so_buffer, LLVMValueRef *wg_offset_dw,
3207 unsigned stream, LLVMValueRef offset_vtx,
3208 LLVMValueRef vertexptr)
3209 {
3210 struct radv_streamout_info *so = &ctx->shader_info->so;
3211 LLVMBuilderRef builder = ctx->ac.builder;
3212 LLVMValueRef offset[4] = {};
3213 LLVMValueRef tmp;
3214
3215 for (unsigned buffer = 0; buffer < 4; ++buffer) {
3216 if (!wg_offset_dw[buffer])
3217 continue;
3218
3219 tmp = LLVMBuildMul(builder, offset_vtx,
3220 LLVMConstInt(ctx->ac.i32, so->strides[buffer], false), "");
3221 tmp = LLVMBuildAdd(builder, wg_offset_dw[buffer], tmp, "");
3222 offset[buffer] = LLVMBuildShl(builder, tmp, LLVMConstInt(ctx->ac.i32, 2, false), "");
3223 }
3224
3225 if (ctx->stage == MESA_SHADER_GEOMETRY) {
3226 struct radv_shader_output_values outputs[AC_LLVM_MAX_OUTPUTS];
3227 unsigned noutput = 0;
3228 unsigned out_idx = 0;
3229
3230 for (unsigned i = 0; i < AC_LLVM_MAX_OUTPUTS; ++i) {
3231 unsigned output_usage_mask =
3232 ctx->shader_info->gs.output_usage_mask[i];
3233 uint8_t output_stream =
3234 output_stream = ctx->shader_info->gs.output_streams[i];
3235
3236 if (!(ctx->output_mask & (1ull << i)) ||
3237 output_stream != stream)
3238 continue;
3239
3240 outputs[noutput].slot_name = i;
3241 outputs[noutput].slot_index = i == VARYING_SLOT_CLIP_DIST1;
3242 outputs[noutput].usage_mask = output_usage_mask;
3243
3244 int length = util_last_bit(output_usage_mask);
3245
3246 for (unsigned j = 0; j < length; j++, out_idx++) {
3247 if (!(output_usage_mask & (1 << j)))
3248 continue;
3249
3250 tmp = ac_build_gep0(&ctx->ac, vertexptr,
3251 LLVMConstInt(ctx->ac.i32, out_idx, false));
3252 outputs[noutput].values[j] = LLVMBuildLoad(builder, tmp, "");
3253 }
3254
3255 for (unsigned j = length; j < 4; j++)
3256 outputs[noutput].values[j] = LLVMGetUndef(ctx->ac.f32);
3257
3258 noutput++;
3259 }
3260
3261 for (unsigned i = 0; i < noutput; i++) {
3262 struct radv_stream_output *output =
3263 radv_get_stream_output_by_loc(so, outputs[i].slot_name);
3264
3265 if (!output ||
3266 output->stream != stream)
3267 continue;
3268
3269 struct radv_shader_output_values out = {};
3270
3271 for (unsigned j = 0; j < 4; j++) {
3272 out.values[j] = outputs[i].values[j];
3273 }
3274
3275 radv_emit_stream_output(ctx, so_buffer, offset, output, &out);
3276 }
3277 } else {
3278 for (unsigned i = 0; i < so->num_outputs; ++i) {
3279 struct radv_stream_output *output =
3280 &ctx->shader_info->so.outputs[i];
3281
3282 if (stream != output->stream)
3283 continue;
3284
3285 struct radv_shader_output_values out = {};
3286
3287 for (unsigned comp = 0; comp < 4; comp++) {
3288 if (!(output->component_mask & (1 << comp)))
3289 continue;
3290
3291 tmp = ac_build_gep0(&ctx->ac, vertexptr,
3292 LLVMConstInt(ctx->ac.i32, 4 * i + comp, false));
3293 out.values[comp] = LLVMBuildLoad(builder, tmp, "");
3294 }
3295
3296 radv_emit_stream_output(ctx, so_buffer, offset, output, &out);
3297 }
3298 }
3299 }
3300
3301 struct ngg_streamout {
3302 LLVMValueRef num_vertices;
3303
3304 /* per-thread data */
3305 LLVMValueRef prim_enable[4]; /* i1 per stream */
3306 LLVMValueRef vertices[3]; /* [N x i32] addrspace(LDS)* */
3307
3308 /* Output */
3309 LLVMValueRef emit[4]; /* per-stream emitted primitives (only valid for used streams) */
3310 };
3311
3312 /**
3313 * Build streamout logic.
3314 *
3315 * Implies a barrier.
3316 *
3317 * Writes number of emitted primitives to gs_ngg_scratch[4:7].
3318 *
3319 * Clobbers gs_ngg_scratch[8:].
3320 */
3321 static void build_streamout(struct radv_shader_context *ctx,
3322 struct ngg_streamout *nggso)
3323 {
3324 struct radv_streamout_info *so = &ctx->shader_info->so;
3325 LLVMBuilderRef builder = ctx->ac.builder;
3326 LLVMValueRef buf_ptr = ctx->streamout_buffers;
3327 LLVMValueRef tid = get_thread_id_in_tg(ctx);
3328 LLVMValueRef cond, tmp, tmp2;
3329 LLVMValueRef i32_2 = LLVMConstInt(ctx->ac.i32, 2, false);
3330 LLVMValueRef i32_4 = LLVMConstInt(ctx->ac.i32, 4, false);
3331 LLVMValueRef i32_8 = LLVMConstInt(ctx->ac.i32, 8, false);
3332 LLVMValueRef so_buffer[4] = {};
3333 unsigned max_num_vertices = 1 + (nggso->vertices[1] ? 1 : 0) +
3334 (nggso->vertices[2] ? 1 : 0);
3335 LLVMValueRef prim_stride_dw[4] = {};
3336 LLVMValueRef prim_stride_dw_vgpr = LLVMGetUndef(ctx->ac.i32);
3337 int stream_for_buffer[4] = { -1, -1, -1, -1 };
3338 unsigned bufmask_for_stream[4] = {};