projects / mesa.git / blob
? search:
summary | shortlog | log | commit | commitdiff | tree
history | raw | HEAD
radv: always load 3 channels for formats that need to be shuffled
[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 "gfx9d.h"
40 #include "ac_binary.h"
41 #include "ac_llvm_util.h"
42 #include "ac_llvm_build.h"
43 #include "ac_shader_abi.h"
44 #include "ac_shader_util.h"
45 #include "ac_exp_param.h"
46
47 #define RADEON_LLVM_MAX_INPUTS (VARYING_SLOT_VAR31 + 1)
48
49 struct radv_shader_context {
50 struct ac_llvm_context ac;
51 const struct radv_nir_compiler_options *options;
52 struct radv_shader_variant_info *shader_info;
53 struct ac_shader_abi abi;
54
55 unsigned max_workgroup_size;
56 LLVMContextRef context;
57 LLVMValueRef main_function;
58
59 LLVMValueRef descriptor_sets[RADV_UD_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 LLVMValueRef gs2vs_offset;
75 LLVMValueRef gs_wave_id;
76 LLVMValueRef gs_vtx_offset[6];
77
78 LLVMValueRef esgs_ring;
79 LLVMValueRef gsvs_ring[4];
80 LLVMValueRef hs_ring_tess_offchip;
81 LLVMValueRef hs_ring_tess_factor;
82
83 LLVMValueRef persp_sample, persp_center, persp_centroid;
84 LLVMValueRef linear_sample, linear_center, linear_centroid;
85
86 /* Streamout */
87 LLVMValueRef streamout_buffers;
88 LLVMValueRef streamout_write_idx;
89 LLVMValueRef streamout_config;
90 LLVMValueRef streamout_offset[4];
91
92 gl_shader_stage stage;
93
94 LLVMValueRef inputs[RADEON_LLVM_MAX_INPUTS * 4];
95 uint64_t float16_shaded_mask;
96
97 uint64_t input_mask;
98 uint64_t output_mask;
99
100 bool is_gs_copy_shader;
101 LLVMValueRef gs_next_vertex[4];
102 unsigned gs_max_out_vertices;
103
104 unsigned tes_primitive_mode;
105
106 uint32_t tcs_patch_outputs_read;
107 uint64_t tcs_outputs_read;
108 uint32_t tcs_vertices_per_patch;
109 uint32_t tcs_num_inputs;
110 uint32_t tcs_num_patches;
111 uint32_t max_gsvs_emit_size;
112 uint32_t gsvs_vertex_size;
113 };
114
115 enum radeon_llvm_calling_convention {
116 RADEON_LLVM_AMDGPU_VS = 87,
117 RADEON_LLVM_AMDGPU_GS = 88,
118 RADEON_LLVM_AMDGPU_PS = 89,
119 RADEON_LLVM_AMDGPU_CS = 90,
120 RADEON_LLVM_AMDGPU_HS = 93,
121 };
122
123 static inline struct radv_shader_context *
124 radv_shader_context_from_abi(struct ac_shader_abi *abi)
125 {
126 struct radv_shader_context *ctx = NULL;
127 return container_of(abi, ctx, abi);
128 }
129
130 struct ac_build_if_state
131 {
132 struct radv_shader_context *ctx;
133 LLVMValueRef condition;
134 LLVMBasicBlockRef entry_block;
135 LLVMBasicBlockRef true_block;
136 LLVMBasicBlockRef false_block;
137 LLVMBasicBlockRef merge_block;
138 };
139
140 static LLVMBasicBlockRef
141 ac_build_insert_new_block(struct radv_shader_context *ctx, const char *name)
142 {
143 LLVMBasicBlockRef current_block;
144 LLVMBasicBlockRef next_block;
145 LLVMBasicBlockRef new_block;
146
147 /* get current basic block */
148 current_block = LLVMGetInsertBlock(ctx->ac.builder);
149
150 /* chqeck if there's another block after this one */
151 next_block = LLVMGetNextBasicBlock(current_block);
152 if (next_block) {
153 /* insert the new block before the next block */
154 new_block = LLVMInsertBasicBlockInContext(ctx->context, next_block, name);
155 }
156 else {
157 /* append new block after current block */
158 LLVMValueRef function = LLVMGetBasicBlockParent(current_block);
159 new_block = LLVMAppendBasicBlockInContext(ctx->context, function, name);
160 }
161 return new_block;
162 }
163
164 static void
165 ac_nir_build_if(struct ac_build_if_state *ifthen,
166 struct radv_shader_context *ctx,
167 LLVMValueRef condition)
168 {
169 LLVMBasicBlockRef block = LLVMGetInsertBlock(ctx->ac.builder);
170
171 memset(ifthen, 0, sizeof *ifthen);
172 ifthen->ctx = ctx;
173 ifthen->condition = condition;
174 ifthen->entry_block = block;
175
176 /* create endif/merge basic block for the phi functions */
177 ifthen->merge_block = ac_build_insert_new_block(ctx, "endif-block");
178
179 /* create/insert true_block before merge_block */
180 ifthen->true_block =
181 LLVMInsertBasicBlockInContext(ctx->context,
182 ifthen->merge_block,
183 "if-true-block");
184
185 /* successive code goes into the true block */
186 LLVMPositionBuilderAtEnd(ctx->ac.builder, ifthen->true_block);
187 }
188
189 /**
190 * End a conditional.
191 */
192 static void
193 ac_nir_build_endif(struct ac_build_if_state *ifthen)
194 {
195 LLVMBuilderRef builder = ifthen->ctx->ac.builder;
196
197 /* Insert branch to the merge block from current block */
198 LLVMBuildBr(builder, ifthen->merge_block);
199
200 /*
201 * Now patch in the various branch instructions.
202 */
203
204 /* Insert the conditional branch instruction at the end of entry_block */
205 LLVMPositionBuilderAtEnd(builder, ifthen->entry_block);
206 if (ifthen->false_block) {
207 /* we have an else clause */
208 LLVMBuildCondBr(builder, ifthen->condition,
209 ifthen->true_block, ifthen->false_block);
210 }
211 else {
212 /* no else clause */
213 LLVMBuildCondBr(builder, ifthen->condition,
214 ifthen->true_block, ifthen->merge_block);
215 }
216
217 /* Resume building code at end of the ifthen->merge_block */
218 LLVMPositionBuilderAtEnd(builder, ifthen->merge_block);
219 }
220
221
222 static LLVMValueRef get_rel_patch_id(struct radv_shader_context *ctx)
223 {
224 switch (ctx->stage) {
225 case MESA_SHADER_TESS_CTRL:
226 return ac_unpack_param(&ctx->ac, ctx->abi.tcs_rel_ids, 0, 8);
227 case MESA_SHADER_TESS_EVAL:
228 return ctx->tes_rel_patch_id;
229 break;
230 default:
231 unreachable("Illegal stage");
232 }
233 }
234
235 static unsigned
236 get_tcs_num_patches(struct radv_shader_context *ctx)
237 {
238 unsigned num_tcs_input_cp = ctx->options->key.tcs.input_vertices;
239 unsigned num_tcs_output_cp = ctx->tcs_vertices_per_patch;
240 uint32_t input_vertex_size = ctx->tcs_num_inputs * 16;
241 uint32_t input_patch_size = ctx->options->key.tcs.input_vertices * input_vertex_size;
242 uint32_t num_tcs_outputs = util_last_bit64(ctx->shader_info->info.tcs.outputs_written);
243 uint32_t num_tcs_patch_outputs = util_last_bit64(ctx->shader_info->info.tcs.patch_outputs_written);
244 uint32_t output_vertex_size = num_tcs_outputs * 16;
245 uint32_t pervertex_output_patch_size = ctx->tcs_vertices_per_patch * output_vertex_size;
246 uint32_t output_patch_size = pervertex_output_patch_size + num_tcs_patch_outputs * 16;
247 unsigned num_patches;
248 unsigned hardware_lds_size;
249
250 /* Ensure that we only need one wave per SIMD so we don't need to check
251 * resource usage. Also ensures that the number of tcs in and out
252 * vertices per threadgroup are at most 256.
253 */
254 num_patches = 64 / MAX2(num_tcs_input_cp, num_tcs_output_cp) * 4;
255 /* Make sure that the data fits in LDS. This assumes the shaders only
256 * use LDS for the inputs and outputs.
257 */
258 hardware_lds_size = 32768;
259
260 /* Looks like STONEY hangs if we use more than 32 KiB LDS in a single
261 * threadgroup, even though there is more than 32 KiB LDS.
262 *
263 * Test: dEQP-VK.tessellation.shader_input_output.barrier
264 */
265 if (ctx->options->chip_class >= CIK && ctx->options->family != CHIP_STONEY)
266 hardware_lds_size = 65536;
267
268 num_patches = MIN2(num_patches, hardware_lds_size / (input_patch_size + output_patch_size));
269 /* Make sure the output data fits in the offchip buffer */
270 num_patches = MIN2(num_patches, (ctx->options->tess_offchip_block_dw_size * 4) / output_patch_size);
271 /* Not necessary for correctness, but improves performance. The
272 * specific value is taken from the proprietary driver.
273 */
274 num_patches = MIN2(num_patches, 40);
275
276 /* SI bug workaround - limit LS-HS threadgroups to only one wave. */
277 if (ctx->options->chip_class == SI) {
278 unsigned one_wave = 64 / MAX2(num_tcs_input_cp, num_tcs_output_cp);
279 num_patches = MIN2(num_patches, one_wave);
280 }
281 return num_patches;
282 }
283
284 static unsigned
285 calculate_tess_lds_size(struct radv_shader_context *ctx)
286 {
287 unsigned num_tcs_input_cp = ctx->options->key.tcs.input_vertices;
288 unsigned num_tcs_output_cp;
289 unsigned num_tcs_outputs, num_tcs_patch_outputs;
290 unsigned input_vertex_size, output_vertex_size;
291 unsigned input_patch_size, output_patch_size;
292 unsigned pervertex_output_patch_size;
293 unsigned output_patch0_offset;
294 unsigned num_patches;
295 unsigned lds_size;
296
297 num_tcs_output_cp = ctx->tcs_vertices_per_patch;
298 num_tcs_outputs = util_last_bit64(ctx->shader_info->info.tcs.outputs_written);
299 num_tcs_patch_outputs = util_last_bit64(ctx->shader_info->info.tcs.patch_outputs_written);
300
301 input_vertex_size = ctx->tcs_num_inputs * 16;
302 output_vertex_size = num_tcs_outputs * 16;
303
304 input_patch_size = num_tcs_input_cp * input_vertex_size;
305
306 pervertex_output_patch_size = num_tcs_output_cp * output_vertex_size;
307 output_patch_size = pervertex_output_patch_size + num_tcs_patch_outputs * 16;
308
309 num_patches = ctx->tcs_num_patches;
310 output_patch0_offset = input_patch_size * num_patches;
311
312 lds_size = output_patch0_offset + output_patch_size * num_patches;
313 return lds_size;
314 }
315
316 /* Tessellation shaders pass outputs to the next shader using LDS.
317 *
318 * LS outputs = TCS inputs
319 * TCS outputs = TES inputs
320 *
321 * The LDS layout is:
322 * - TCS inputs for patch 0
323 * - TCS inputs for patch 1
324 * - TCS inputs for patch 2 = get_tcs_in_current_patch_offset (if RelPatchID==2)
325 * - ...
326 * - TCS outputs for patch 0 = get_tcs_out_patch0_offset
327 * - Per-patch TCS outputs for patch 0 = get_tcs_out_patch0_patch_data_offset
328 * - TCS outputs for patch 1
329 * - Per-patch TCS outputs for patch 1
330 * - TCS outputs for patch 2 = get_tcs_out_current_patch_offset (if RelPatchID==2)
331 * - Per-patch TCS outputs for patch 2 = get_tcs_out_current_patch_data_offset (if RelPatchID==2)
332 * - ...
333 *
334 * All three shaders VS(LS), TCS, TES share the same LDS space.
335 */
336 static LLVMValueRef
337 get_tcs_in_patch_stride(struct radv_shader_context *ctx)
338 {
339 assert (ctx->stage == MESA_SHADER_TESS_CTRL);
340 uint32_t input_vertex_size = ctx->tcs_num_inputs * 16;
341 uint32_t input_patch_size = ctx->options->key.tcs.input_vertices * input_vertex_size;
342
343 input_patch_size /= 4;
344 return LLVMConstInt(ctx->ac.i32, input_patch_size, false);
345 }
346
347 static LLVMValueRef
348 get_tcs_out_patch_stride(struct radv_shader_context *ctx)
349 {
350 uint32_t num_tcs_outputs = util_last_bit64(ctx->shader_info->info.tcs.outputs_written);
351 uint32_t num_tcs_patch_outputs = util_last_bit64(ctx->shader_info->info.tcs.patch_outputs_written);
352 uint32_t output_vertex_size = num_tcs_outputs * 16;
353 uint32_t pervertex_output_patch_size = ctx->tcs_vertices_per_patch * output_vertex_size;
354 uint32_t output_patch_size = pervertex_output_patch_size + num_tcs_patch_outputs * 16;
355 output_patch_size /= 4;
356 return LLVMConstInt(ctx->ac.i32, output_patch_size, false);
357 }
358
359 static LLVMValueRef
360 get_tcs_out_vertex_stride(struct radv_shader_context *ctx)
361 {
362 uint32_t num_tcs_outputs = util_last_bit64(ctx->shader_info->info.tcs.outputs_written);
363 uint32_t output_vertex_size = num_tcs_outputs * 16;
364 output_vertex_size /= 4;
365 return LLVMConstInt(ctx->ac.i32, output_vertex_size, false);
366 }
367
368 static LLVMValueRef
369 get_tcs_out_patch0_offset(struct radv_shader_context *ctx)
370 {
371 assert (ctx->stage == MESA_SHADER_TESS_CTRL);
372 uint32_t input_vertex_size = ctx->tcs_num_inputs * 16;
373 uint32_t input_patch_size = ctx->options->key.tcs.input_vertices * input_vertex_size;
374 uint32_t output_patch0_offset = input_patch_size;
375 unsigned num_patches = ctx->tcs_num_patches;
376
377 output_patch0_offset *= num_patches;
378 output_patch0_offset /= 4;
379 return LLVMConstInt(ctx->ac.i32, output_patch0_offset, false);
380 }
381
382 static LLVMValueRef
383 get_tcs_out_patch0_patch_data_offset(struct radv_shader_context *ctx)
384 {
385 assert (ctx->stage == MESA_SHADER_TESS_CTRL);
386 uint32_t input_vertex_size = ctx->tcs_num_inputs * 16;
387 uint32_t input_patch_size = ctx->options->key.tcs.input_vertices * input_vertex_size;
388 uint32_t output_patch0_offset = input_patch_size;
389
390 uint32_t num_tcs_outputs = util_last_bit64(ctx->shader_info->info.tcs.outputs_written);
391 uint32_t output_vertex_size = num_tcs_outputs * 16;
392 uint32_t pervertex_output_patch_size = ctx->tcs_vertices_per_patch * output_vertex_size;
393 unsigned num_patches = ctx->tcs_num_patches;
394
395 output_patch0_offset *= num_patches;
396 output_patch0_offset += pervertex_output_patch_size;
397 output_patch0_offset /= 4;
398 return LLVMConstInt(ctx->ac.i32, output_patch0_offset, false);
399 }
400
401 static LLVMValueRef
402 get_tcs_in_current_patch_offset(struct radv_shader_context *ctx)
403 {
404 LLVMValueRef patch_stride = get_tcs_in_patch_stride(ctx);
405 LLVMValueRef rel_patch_id = get_rel_patch_id(ctx);
406
407 return LLVMBuildMul(ctx->ac.builder, patch_stride, rel_patch_id, "");
408 }
409
410 static LLVMValueRef
411 get_tcs_out_current_patch_offset(struct radv_shader_context *ctx)
412 {
413 LLVMValueRef patch0_offset = get_tcs_out_patch0_offset(ctx);
414 LLVMValueRef patch_stride = get_tcs_out_patch_stride(ctx);
415 LLVMValueRef rel_patch_id = get_rel_patch_id(ctx);
416
417 return ac_build_imad(&ctx->ac, patch_stride, rel_patch_id,
418 patch0_offset);
419 }
420
421 static LLVMValueRef
422 get_tcs_out_current_patch_data_offset(struct radv_shader_context *ctx)
423 {
424 LLVMValueRef patch0_patch_data_offset =
425 get_tcs_out_patch0_patch_data_offset(ctx);
426 LLVMValueRef patch_stride = get_tcs_out_patch_stride(ctx);
427 LLVMValueRef rel_patch_id = get_rel_patch_id(ctx);
428
429 return ac_build_imad(&ctx->ac, patch_stride, rel_patch_id,
430 patch0_patch_data_offset);
431 }
432
433 #define MAX_ARGS 64
434 struct arg_info {
435 LLVMTypeRef types[MAX_ARGS];
436 LLVMValueRef *assign[MAX_ARGS];
437 uint8_t count;
438 uint8_t sgpr_count;
439 uint8_t num_sgprs_used;
440 uint8_t num_vgprs_used;
441 };
442
443 enum ac_arg_regfile {
444 ARG_SGPR,
445 ARG_VGPR,
446 };
447
448 static void
449 add_arg(struct arg_info *info, enum ac_arg_regfile regfile, LLVMTypeRef type,
450 LLVMValueRef *param_ptr)
451 {
452 assert(info->count < MAX_ARGS);
453
454 info->assign[info->count] = param_ptr;
455 info->types[info->count] = type;
456 info->count++;
457
458 if (regfile == ARG_SGPR) {
459 info->num_sgprs_used += ac_get_type_size(type) / 4;
460 info->sgpr_count++;
461 } else {
462 assert(regfile == ARG_VGPR);
463 info->num_vgprs_used += ac_get_type_size(type) / 4;
464 }
465 }
466
467 static void assign_arguments(LLVMValueRef main_function,
468 struct arg_info *info)
469 {
470 unsigned i;
471 for (i = 0; i < info->count; i++) {
472 if (info->assign[i])
473 *info->assign[i] = LLVMGetParam(main_function, i);
474 }
475 }
476
477 static LLVMValueRef
478 create_llvm_function(LLVMContextRef ctx, LLVMModuleRef module,
479 LLVMBuilderRef builder, LLVMTypeRef *return_types,
480 unsigned num_return_elems,
481 struct arg_info *args,
482 unsigned max_workgroup_size,
483 const struct radv_nir_compiler_options *options)
484 {
485 LLVMTypeRef main_function_type, ret_type;
486 LLVMBasicBlockRef main_function_body;
487
488 if (num_return_elems)
489 ret_type = LLVMStructTypeInContext(ctx, return_types,
490 num_return_elems, true);
491 else
492 ret_type = LLVMVoidTypeInContext(ctx);
493
494 /* Setup the function */
495 main_function_type =
496 LLVMFunctionType(ret_type, args->types, args->count, 0);
497 LLVMValueRef main_function =
498 LLVMAddFunction(module, "main", main_function_type);
499 main_function_body =
500 LLVMAppendBasicBlockInContext(ctx, main_function, "main_body");
501 LLVMPositionBuilderAtEnd(builder, main_function_body);
502
503 LLVMSetFunctionCallConv(main_function, RADEON_LLVM_AMDGPU_CS);
504 for (unsigned i = 0; i < args->sgpr_count; ++i) {
505 LLVMValueRef P = LLVMGetParam(main_function, i);
506
507 ac_add_function_attr(ctx, main_function, i + 1, AC_FUNC_ATTR_INREG);
508
509 if (LLVMGetTypeKind(LLVMTypeOf(P)) == LLVMPointerTypeKind) {
510 ac_add_function_attr(ctx, main_function, i + 1, AC_FUNC_ATTR_NOALIAS);
511 ac_add_attr_dereferenceable(P, UINT64_MAX);
512 }
513 }
514
515 if (options->address32_hi) {
516 ac_llvm_add_target_dep_function_attr(main_function,
517 "amdgpu-32bit-address-high-bits",
518 options->address32_hi);
519 }
520
521 if (max_workgroup_size) {
522 ac_llvm_add_target_dep_function_attr(main_function,
523 "amdgpu-max-work-group-size",
524 max_workgroup_size);
525 }
526 if (options->unsafe_math) {
527 /* These were copied from some LLVM test. */
528 LLVMAddTargetDependentFunctionAttr(main_function,
529 "less-precise-fpmad",
530 "true");
531 LLVMAddTargetDependentFunctionAttr(main_function,
532 "no-infs-fp-math",
533 "true");
534 LLVMAddTargetDependentFunctionAttr(main_function,
535 "no-nans-fp-math",
536 "true");
537 LLVMAddTargetDependentFunctionAttr(main_function,
538 "unsafe-fp-math",
539 "true");
540 LLVMAddTargetDependentFunctionAttr(main_function,
541 "no-signed-zeros-fp-math",
542 "true");
543 }
544 return main_function;
545 }
546
547
548 static void
549 set_loc(struct radv_userdata_info *ud_info, uint8_t *sgpr_idx,
550 uint8_t num_sgprs)
551 {
552 ud_info->sgpr_idx = *sgpr_idx;
553 ud_info->num_sgprs = num_sgprs;
554 *sgpr_idx += num_sgprs;
555 }
556
557 static void
558 set_loc_shader(struct radv_shader_context *ctx, int idx, uint8_t *sgpr_idx,
559 uint8_t num_sgprs)
560 {
561 struct radv_userdata_info *ud_info =
562 &ctx->shader_info->user_sgprs_locs.shader_data[idx];
563 assert(ud_info);
564
565 set_loc(ud_info, sgpr_idx, num_sgprs);
566 }
567
568 static void
569 set_loc_shader_ptr(struct radv_shader_context *ctx, int idx, uint8_t *sgpr_idx)
570 {
571 bool use_32bit_pointers = idx != AC_UD_SCRATCH_RING_OFFSETS;
572
573 set_loc_shader(ctx, idx, sgpr_idx, use_32bit_pointers ? 1 : 2);
574 }
575
576 static void
577 set_loc_desc(struct radv_shader_context *ctx, int idx, uint8_t *sgpr_idx)
578 {
579 struct radv_userdata_locations *locs =
580 &ctx->shader_info->user_sgprs_locs;
581 struct radv_userdata_info *ud_info = &locs->descriptor_sets[idx];
582 assert(ud_info);
583
584 set_loc(ud_info, sgpr_idx, 1);
585
586 locs->descriptor_sets_enabled |= 1 << idx;
587 }
588
589 struct user_sgpr_info {
590 bool need_ring_offsets;
591 bool indirect_all_descriptor_sets;
592 uint8_t remaining_sgprs;
593 };
594
595 static bool needs_view_index_sgpr(struct radv_shader_context *ctx,
596 gl_shader_stage stage)
597 {
598 switch (stage) {
599 case MESA_SHADER_VERTEX:
600 if (ctx->shader_info->info.needs_multiview_view_index ||
601 (!ctx->options->key.vs.as_es && !ctx->options->key.vs.as_ls && ctx->options->key.has_multiview_view_index))
602 return true;
603 break;
604 case MESA_SHADER_TESS_EVAL:
605 if (ctx->shader_info->info.needs_multiview_view_index || (!ctx->options->key.tes.as_es && ctx->options->key.has_multiview_view_index))
606 return true;
607 break;
608 case MESA_SHADER_GEOMETRY:
609 case MESA_SHADER_TESS_CTRL:
610 if (ctx->shader_info->info.needs_multiview_view_index)
611 return true;
612 break;
613 default:
614 break;
615 }
616 return false;
617 }
618
619 static uint8_t
620 count_vs_user_sgprs(struct radv_shader_context *ctx)
621 {
622 uint8_t count = 0;
623
624 if (ctx->shader_info->info.vs.has_vertex_buffers)
625 count++;
626 count += ctx->shader_info->info.vs.needs_draw_id ? 3 : 2;
627
628 return count;
629 }
630
631 static void allocate_inline_push_consts(struct radv_shader_context *ctx,
632 struct user_sgpr_info *user_sgpr_info)
633 {
634 uint8_t remaining_sgprs = user_sgpr_info->remaining_sgprs;
635
636 /* Only supported if shaders use push constants. */
637 if (ctx->shader_info->info.min_push_constant_used == UINT8_MAX)
638 return;
639
640 /* Only supported if shaders don't have indirect push constants. */
641 if (ctx->shader_info->info.has_indirect_push_constants)
642 return;
643
644 /* Only supported for 32-bit push constants. */
645 if (!ctx->shader_info->info.has_only_32bit_push_constants)
646 return;
647
648 uint8_t num_push_consts =
649 (ctx->shader_info->info.max_push_constant_used -
650 ctx->shader_info->info.min_push_constant_used) / 4;
651
652 /* Check if the number of user SGPRs is large enough. */
653 if (num_push_consts < remaining_sgprs) {
654 ctx->shader_info->info.num_inline_push_consts = num_push_consts;
655 } else {
656 ctx->shader_info->info.num_inline_push_consts = remaining_sgprs;
657 }
658
659 /* Clamp to the maximum number of allowed inlined push constants. */
660 if (ctx->shader_info->info.num_inline_push_consts > AC_MAX_INLINE_PUSH_CONSTS)
661 ctx->shader_info->info.num_inline_push_consts = AC_MAX_INLINE_PUSH_CONSTS;
662
663 if (ctx->shader_info->info.num_inline_push_consts == num_push_consts &&
664 !ctx->shader_info->info.loads_dynamic_offsets) {
665 /* Disable the default push constants path if all constants are
666 * inlined and if shaders don't use dynamic descriptors.
667 */
668 ctx->shader_info->info.loads_push_constants = false;
669 }
670
671 ctx->shader_info->info.base_inline_push_consts =
672 ctx->shader_info->info.min_push_constant_used / 4;
673 }
674
675 static void allocate_user_sgprs(struct radv_shader_context *ctx,
676 gl_shader_stage stage,
677 bool has_previous_stage,
678 gl_shader_stage previous_stage,
679 bool needs_view_index,
680 struct user_sgpr_info *user_sgpr_info)
681 {
682 uint8_t user_sgpr_count = 0;
683
684 memset(user_sgpr_info, 0, sizeof(struct user_sgpr_info));
685
686 /* until we sort out scratch/global buffers always assign ring offsets for gs/vs/es */
687 if (stage == MESA_SHADER_GEOMETRY ||
688 stage == MESA_SHADER_VERTEX ||
689 stage == MESA_SHADER_TESS_CTRL ||
690 stage == MESA_SHADER_TESS_EVAL ||
691 ctx->is_gs_copy_shader)
692 user_sgpr_info->need_ring_offsets = true;
693
694 if (stage == MESA_SHADER_FRAGMENT &&
695 ctx->shader_info->info.ps.needs_sample_positions)
696 user_sgpr_info->need_ring_offsets = true;
697
698 /* 2 user sgprs will nearly always be allocated for scratch/rings */
699 if (ctx->options->supports_spill || user_sgpr_info->need_ring_offsets) {
700 user_sgpr_count += 2;
701 }
702
703 switch (stage) {
704 case MESA_SHADER_COMPUTE:
705 if (ctx->shader_info->info.cs.uses_grid_size)
706 user_sgpr_count += 3;
707 break;
708 case MESA_SHADER_FRAGMENT:
709 user_sgpr_count += ctx->shader_info->info.ps.needs_sample_positions;
710 break;
711 case MESA_SHADER_VERTEX:
712 if (!ctx->is_gs_copy_shader)
713 user_sgpr_count += count_vs_user_sgprs(ctx);
714 break;
715 case MESA_SHADER_TESS_CTRL:
716 if (has_previous_stage) {
717 if (previous_stage == MESA_SHADER_VERTEX)
718 user_sgpr_count += count_vs_user_sgprs(ctx);
719 }
720 break;
721 case MESA_SHADER_TESS_EVAL:
722 break;
723 case MESA_SHADER_GEOMETRY:
724 if (has_previous_stage) {
725 if (previous_stage == MESA_SHADER_VERTEX) {
726 user_sgpr_count += count_vs_user_sgprs(ctx);
727 }
728 }
729 break;
730 default:
731 break;
732 }
733
734 if (needs_view_index)
735 user_sgpr_count++;
736
737 if (ctx->shader_info->info.loads_push_constants)
738 user_sgpr_count++;
739
740 if (ctx->streamout_buffers)
741 user_sgpr_count++;
742
743 uint32_t available_sgprs = ctx->options->chip_class >= GFX9 && stage != MESA_SHADER_COMPUTE ? 32 : 16;
744 uint32_t remaining_sgprs = available_sgprs - user_sgpr_count;
745 uint32_t num_desc_set =
746 util_bitcount(ctx->shader_info->info.desc_set_used_mask);
747
748 if (remaining_sgprs < num_desc_set) {
749 user_sgpr_info->indirect_all_descriptor_sets = true;
750 user_sgpr_info->remaining_sgprs = remaining_sgprs - 1;
751 } else {
752 user_sgpr_info->remaining_sgprs = remaining_sgprs - num_desc_set;
753 }
754
755 allocate_inline_push_consts(ctx, user_sgpr_info);
756 }
757
758 static void
759 declare_global_input_sgprs(struct radv_shader_context *ctx,
760 const struct user_sgpr_info *user_sgpr_info,
761 struct arg_info *args,
762 LLVMValueRef *desc_sets)
763 {
764 LLVMTypeRef type = ac_array_in_const32_addr_space(ctx->ac.i8);
765
766 /* 1 for each descriptor set */
767 if (!user_sgpr_info->indirect_all_descriptor_sets) {
768 uint32_t mask = ctx->shader_info->info.desc_set_used_mask;
769
770 while (mask) {
771 int i = u_bit_scan(&mask);
772
773 add_arg(args, ARG_SGPR, type, &ctx->descriptor_sets[i]);
774 }
775 } else {
776 add_arg(args, ARG_SGPR, ac_array_in_const32_addr_space(type),
777 desc_sets);
778 }
779
780 if (ctx->shader_info->info.loads_push_constants) {
781 /* 1 for push constants and dynamic descriptors */
782 add_arg(args, ARG_SGPR, type, &ctx->abi.push_constants);
783 }
784
785 for (unsigned i = 0; i < ctx->shader_info->info.num_inline_push_consts; i++) {
786 add_arg(args, ARG_SGPR, ctx->ac.i32,
787 &ctx->abi.inline_push_consts[i]);
788 }
789 ctx->abi.num_inline_push_consts = ctx->shader_info->info.num_inline_push_consts;
790 ctx->abi.base_inline_push_consts = ctx->shader_info->info.base_inline_push_consts;
791
792 if (ctx->shader_info->info.so.num_outputs) {
793 add_arg(args, ARG_SGPR,
794 ac_array_in_const32_addr_space(ctx->ac.v4i32),
795 &ctx->streamout_buffers);
796 }
797 }
798
799 static void
800 declare_vs_specific_input_sgprs(struct radv_shader_context *ctx,
801 gl_shader_stage stage,
802 bool has_previous_stage,
803 gl_shader_stage previous_stage,
804 struct arg_info *args)
805 {
806 if (!ctx->is_gs_copy_shader &&
807 (stage == MESA_SHADER_VERTEX ||
808 (has_previous_stage && previous_stage == MESA_SHADER_VERTEX))) {
809 if (ctx->shader_info->info.vs.has_vertex_buffers) {
810 add_arg(args, ARG_SGPR,
811 ac_array_in_const32_addr_space(ctx->ac.v4i32),
812 &ctx->vertex_buffers);
813 }
814 add_arg(args, ARG_SGPR, ctx->ac.i32, &ctx->abi.base_vertex);
815 add_arg(args, ARG_SGPR, ctx->ac.i32, &ctx->abi.start_instance);
816 if (ctx->shader_info->info.vs.needs_draw_id) {
817 add_arg(args, ARG_SGPR, ctx->ac.i32, &ctx->abi.draw_id);
818 }
819 }
820 }
821
822 static void
823 declare_vs_input_vgprs(struct radv_shader_context *ctx, struct arg_info *args)
824 {
825 add_arg(args, ARG_VGPR, ctx->ac.i32, &ctx->abi.vertex_id);
826 if (!ctx->is_gs_copy_shader) {
827 if (ctx->options->key.vs.as_ls) {
828 add_arg(args, ARG_VGPR, ctx->ac.i32, &ctx->rel_auto_id);
829 add_arg(args, ARG_VGPR, ctx->ac.i32, &ctx->abi.instance_id);
830 } else {
831 add_arg(args, ARG_VGPR, ctx->ac.i32, &ctx->abi.instance_id);
832 add_arg(args, ARG_VGPR, ctx->ac.i32, &ctx->vs_prim_id);
833 }
834 add_arg(args, ARG_VGPR, ctx->ac.i32, NULL); /* unused */
835 }
836 }
837
838 static void
839 declare_streamout_sgprs(struct radv_shader_context *ctx, gl_shader_stage stage,
840 struct arg_info *args)
841 {
842 int i;
843
844 /* Streamout SGPRs. */
845 if (ctx->shader_info->info.so.num_outputs) {
846 assert(stage == MESA_SHADER_VERTEX ||
847 stage == MESA_SHADER_TESS_EVAL);
848
849 if (stage != MESA_SHADER_TESS_EVAL) {
850 add_arg(args, ARG_SGPR, ctx->ac.i32, &ctx->streamout_config);
851 } else {
852 args->assign[args->count - 1] = &ctx->streamout_config;
853 args->types[args->count - 1] = ctx->ac.i32;
854 }
855
856 add_arg(args, ARG_SGPR, ctx->ac.i32, &ctx->streamout_write_idx);
857 }
858
859 /* A streamout buffer offset is loaded if the stride is non-zero. */
860 for (i = 0; i < 4; i++) {
861 if (!ctx->shader_info->info.so.strides[i])
862 continue;
863
864 add_arg(args, ARG_SGPR, ctx->ac.i32, &ctx->streamout_offset[i]);
865 }
866 }
867
868 static void
869 declare_tes_input_vgprs(struct radv_shader_context *ctx, struct arg_info *args)
870 {
871 add_arg(args, ARG_VGPR, ctx->ac.f32, &ctx->tes_u);
872 add_arg(args, ARG_VGPR, ctx->ac.f32, &ctx->tes_v);
873 add_arg(args, ARG_VGPR, ctx->ac.i32, &ctx->tes_rel_patch_id);
874 add_arg(args, ARG_VGPR, ctx->ac.i32, &ctx->abi.tes_patch_id);
875 }
876
877 static void
878 set_global_input_locs(struct radv_shader_context *ctx,
879 const struct user_sgpr_info *user_sgpr_info,
880 LLVMValueRef desc_sets, uint8_t *user_sgpr_idx)
881 {
882 uint32_t mask = ctx->shader_info->info.desc_set_used_mask;
883
884 if (!user_sgpr_info->indirect_all_descriptor_sets) {
885 while (mask) {
886 int i = u_bit_scan(&mask);
887
888 set_loc_desc(ctx, i, user_sgpr_idx);
889 }
890 } else {
891 set_loc_shader_ptr(ctx, AC_UD_INDIRECT_DESCRIPTOR_SETS,
892 user_sgpr_idx);
893
894 while (mask) {
895 int i = u_bit_scan(&mask);
896
897 ctx->descriptor_sets[i] =
898 ac_build_load_to_sgpr(&ctx->ac, desc_sets,
899 LLVMConstInt(ctx->ac.i32, i, false));
900
901 }
902
903 ctx->shader_info->need_indirect_descriptor_sets = true;
904 }
905
906 if (ctx->shader_info->info.loads_push_constants) {
907 set_loc_shader_ptr(ctx, AC_UD_PUSH_CONSTANTS, user_sgpr_idx);
908 }
909
910 if (ctx->shader_info->info.num_inline_push_consts) {
911 set_loc_shader(ctx, AC_UD_INLINE_PUSH_CONSTANTS, user_sgpr_idx,
912 ctx->shader_info->info.num_inline_push_consts);
913 }
914
915 if (ctx->streamout_buffers) {
916 set_loc_shader_ptr(ctx, AC_UD_STREAMOUT_BUFFERS,
917 user_sgpr_idx);
918 }
919 }
920
921 static void
922 set_vs_specific_input_locs(struct radv_shader_context *ctx,
923 gl_shader_stage stage, bool has_previous_stage,
924 gl_shader_stage previous_stage,
925 uint8_t *user_sgpr_idx)
926 {
927 if (!ctx->is_gs_copy_shader &&
928 (stage == MESA_SHADER_VERTEX ||
929 (has_previous_stage && previous_stage == MESA_SHADER_VERTEX))) {
930 if (ctx->shader_info->info.vs.has_vertex_buffers) {
931 set_loc_shader_ptr(ctx, AC_UD_VS_VERTEX_BUFFERS,
932 user_sgpr_idx);
933 }
934
935 unsigned vs_num = 2;
936 if (ctx->shader_info->info.vs.needs_draw_id)
937 vs_num++;
938
939 set_loc_shader(ctx, AC_UD_VS_BASE_VERTEX_START_INSTANCE,
940 user_sgpr_idx, vs_num);
941 }
942 }
943
944 static void set_llvm_calling_convention(LLVMValueRef func,
945 gl_shader_stage stage)
946 {
947 enum radeon_llvm_calling_convention calling_conv;
948
949 switch (stage) {
950 case MESA_SHADER_VERTEX:
951 case MESA_SHADER_TESS_EVAL:
952 calling_conv = RADEON_LLVM_AMDGPU_VS;
953 break;
954 case MESA_SHADER_GEOMETRY:
955 calling_conv = RADEON_LLVM_AMDGPU_GS;
956 break;
957 case MESA_SHADER_TESS_CTRL:
958 calling_conv = RADEON_LLVM_AMDGPU_HS;
959 break;
960 case MESA_SHADER_FRAGMENT:
961 calling_conv = RADEON_LLVM_AMDGPU_PS;
962 break;
963 case MESA_SHADER_COMPUTE:
964 calling_conv = RADEON_LLVM_AMDGPU_CS;
965 break;
966 default:
967 unreachable("Unhandle shader type");
968 }
969
970 LLVMSetFunctionCallConv(func, calling_conv);
971 }
972
973 static void create_function(struct radv_shader_context *ctx,
974 gl_shader_stage stage,
975 bool has_previous_stage,
976 gl_shader_stage previous_stage)
977 {
978 uint8_t user_sgpr_idx;
979 struct user_sgpr_info user_sgpr_info;
980 struct arg_info args = {};
981 LLVMValueRef desc_sets;
982 bool needs_view_index = needs_view_index_sgpr(ctx, stage);
983 allocate_user_sgprs(ctx, stage, has_previous_stage,
984 previous_stage, needs_view_index, &user_sgpr_info);
985
986 if (user_sgpr_info.need_ring_offsets && !ctx->options->supports_spill) {
987 add_arg(&args, ARG_SGPR, ac_array_in_const_addr_space(ctx->ac.v4i32),
988 &ctx->ring_offsets);
989 }
990
991 switch (stage) {
992 case MESA_SHADER_COMPUTE:
993 declare_global_input_sgprs(ctx, &user_sgpr_info, &args,
994 &desc_sets);
995
996 if (ctx->shader_info->info.cs.uses_grid_size) {
997 add_arg(&args, ARG_SGPR, ctx->ac.v3i32,
998 &ctx->abi.num_work_groups);
999 }
1000
1001 for (int i = 0; i < 3; i++) {
1002 ctx->abi.workgroup_ids[i] = NULL;
1003 if (ctx->shader_info->info.cs.uses_block_id[i]) {
1004 add_arg(&args, ARG_SGPR, ctx->ac.i32,
1005 &ctx->abi.workgroup_ids[i]);
1006 }
1007 }
1008
1009 if (ctx->shader_info->info.cs.uses_local_invocation_idx)
1010 add_arg(&args, ARG_SGPR, ctx->ac.i32, &ctx->abi.tg_size);
1011 add_arg(&args, ARG_VGPR, ctx->ac.v3i32,
1012 &ctx->abi.local_invocation_ids);
1013 break;
1014 case MESA_SHADER_VERTEX:
1015 declare_global_input_sgprs(ctx, &user_sgpr_info, &args,
1016 &desc_sets);
1017
1018 declare_vs_specific_input_sgprs(ctx, stage, has_previous_stage,
1019 previous_stage, &args);
1020
1021 if (needs_view_index)
1022 add_arg(&args, ARG_SGPR, ctx->ac.i32,
1023 &ctx->abi.view_index);
1024 if (ctx->options->key.vs.as_es) {
1025 add_arg(&args, ARG_SGPR, ctx->ac.i32,
1026 &ctx->es2gs_offset);
1027 } else if (ctx->options->key.vs.as_ls) {
1028 /* no extra parameters */
1029 } else {
1030 declare_streamout_sgprs(ctx, stage, &args);
1031 }
1032
1033 declare_vs_input_vgprs(ctx, &args);
1034 break;
1035 case MESA_SHADER_TESS_CTRL:
1036 if (has_previous_stage) {
1037 // First 6 system regs
1038 add_arg(&args, ARG_SGPR, ctx->ac.i32, &ctx->oc_lds);
1039 add_arg(&args, ARG_SGPR, ctx->ac.i32,
1040 &ctx->merged_wave_info);
1041 add_arg(&args, ARG_SGPR, ctx->ac.i32,
1042 &ctx->tess_factor_offset);
1043
1044 add_arg(&args, ARG_SGPR, ctx->ac.i32, NULL); // scratch offset
1045 add_arg(&args, ARG_SGPR, ctx->ac.i32, NULL); // unknown
1046 add_arg(&args, ARG_SGPR, ctx->ac.i32, NULL); // unknown
1047
1048 declare_global_input_sgprs(ctx, &user_sgpr_info, &args,
1049 &desc_sets);
1050
1051 declare_vs_specific_input_sgprs(ctx, stage,
1052 has_previous_stage,
1053 previous_stage, &args);
1054
1055 if (needs_view_index)
1056 add_arg(&args, ARG_SGPR, ctx->ac.i32,
1057 &ctx->abi.view_index);
1058
1059 add_arg(&args, ARG_VGPR, ctx->ac.i32,
1060 &ctx->abi.tcs_patch_id);
1061 add_arg(&args, ARG_VGPR, ctx->ac.i32,
1062 &ctx->abi.tcs_rel_ids);
1063
1064 declare_vs_input_vgprs(ctx, &args);
1065 } else {
1066 declare_global_input_sgprs(ctx, &user_sgpr_info, &args,
1067 &desc_sets);
1068
1069 if (needs_view_index)
1070 add_arg(&args, ARG_SGPR, ctx->ac.i32,
1071 &ctx->abi.view_index);
1072
1073 add_arg(&args, ARG_SGPR, ctx->ac.i32, &ctx->oc_lds);
1074 add_arg(&args, ARG_SGPR, ctx->ac.i32,
1075 &ctx->tess_factor_offset);
1076 add_arg(&args, ARG_VGPR, ctx->ac.i32,
1077 &ctx->abi.tcs_patch_id);
1078 add_arg(&args, ARG_VGPR, ctx->ac.i32,
1079 &ctx->abi.tcs_rel_ids);
1080 }
1081 break;
1082 case MESA_SHADER_TESS_EVAL:
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 if (ctx->options->key.tes.as_es) {
1091 add_arg(&args, ARG_SGPR, ctx->ac.i32, &ctx->oc_lds);
1092 add_arg(&args, ARG_SGPR, ctx->ac.i32, NULL);
1093 add_arg(&args, ARG_SGPR, ctx->ac.i32,
1094 &ctx->es2gs_offset);
1095 } else {
1096 add_arg(&args, ARG_SGPR, ctx->ac.i32, NULL);
1097 declare_streamout_sgprs(ctx, stage, &args);
1098 add_arg(&args, ARG_SGPR, ctx->ac.i32, &ctx->oc_lds);
1099 }
1100 declare_tes_input_vgprs(ctx, &args);
1101 break;
1102 case MESA_SHADER_GEOMETRY:
1103 if (has_previous_stage) {
1104 // First 6 system regs
1105 add_arg(&args, ARG_SGPR, ctx->ac.i32,
1106 &ctx->gs2vs_offset);
1107 add_arg(&args, ARG_SGPR, ctx->ac.i32,
1108 &ctx->merged_wave_info);
1109 add_arg(&args, ARG_SGPR, ctx->ac.i32, &ctx->oc_lds);
1110
1111 add_arg(&args, ARG_SGPR, ctx->ac.i32, NULL); // scratch offset
1112 add_arg(&args, ARG_SGPR, ctx->ac.i32, NULL); // unknown
1113 add_arg(&args, ARG_SGPR, ctx->ac.i32, NULL); // unknown
1114
1115 declare_global_input_sgprs(ctx, &user_sgpr_info, &args,
1116 &desc_sets);
1117
1118 if (previous_stage != MESA_SHADER_TESS_EVAL) {
1119 declare_vs_specific_input_sgprs(ctx, stage,
1120 has_previous_stage,
1121 previous_stage,
1122 &args);
1123 }
1124
1125 if (needs_view_index)
1126 add_arg(&args, ARG_SGPR, ctx->ac.i32,
1127 &ctx->abi.view_index);
1128
1129 add_arg(&args, ARG_VGPR, ctx->ac.i32,
1130 &ctx->gs_vtx_offset[0]);
1131 add_arg(&args, ARG_VGPR, ctx->ac.i32,
1132 &ctx->gs_vtx_offset[2]);
1133 add_arg(&args, ARG_VGPR, ctx->ac.i32,
1134 &ctx->abi.gs_prim_id);
1135 add_arg(&args, ARG_VGPR, ctx->ac.i32,
1136 &ctx->abi.gs_invocation_id);
1137 add_arg(&args, ARG_VGPR, ctx->ac.i32,
1138 &ctx->gs_vtx_offset[4]);
1139
1140 if (previous_stage == MESA_SHADER_VERTEX) {
1141 declare_vs_input_vgprs(ctx, &args);
1142 } else {
1143 declare_tes_input_vgprs(ctx, &args);
1144 }
1145 } else {
1146 declare_global_input_sgprs(ctx, &user_sgpr_info, &args,
1147 &desc_sets);
1148
1149 if (needs_view_index)
1150 add_arg(&args, ARG_SGPR, ctx->ac.i32,
1151 &ctx->abi.view_index);
1152
1153 add_arg(&args, ARG_SGPR, ctx->ac.i32, &ctx->gs2vs_offset);
1154 add_arg(&args, ARG_SGPR, ctx->ac.i32, &ctx->gs_wave_id);
1155 add_arg(&args, ARG_VGPR, ctx->ac.i32,
1156 &ctx->gs_vtx_offset[0]);
1157 add_arg(&args, ARG_VGPR, ctx->ac.i32,
1158 &ctx->gs_vtx_offset[1]);
1159 add_arg(&args, ARG_VGPR, ctx->ac.i32,
1160 &ctx->abi.gs_prim_id);
1161 add_arg(&args, ARG_VGPR, ctx->ac.i32,
1162 &ctx->gs_vtx_offset[2]);
1163 add_arg(&args, ARG_VGPR, ctx->ac.i32,
1164 &ctx->gs_vtx_offset[3]);
1165 add_arg(&args, ARG_VGPR, ctx->ac.i32,
1166 &ctx->gs_vtx_offset[4]);
1167 add_arg(&args, ARG_VGPR, ctx->ac.i32,
1168 &ctx->gs_vtx_offset[5]);
1169 add_arg(&args, ARG_VGPR, ctx->ac.i32,
1170 &ctx->abi.gs_invocation_id);
1171 }
1172 break;
1173 case MESA_SHADER_FRAGMENT:
1174 declare_global_input_sgprs(ctx, &user_sgpr_info, &args,
1175 &desc_sets);
1176
1177 add_arg(&args, ARG_SGPR, ctx->ac.i32, &ctx->abi.prim_mask);
1178 add_arg(&args, ARG_VGPR, ctx->ac.v2i32, &ctx->persp_sample);
1179 add_arg(&args, ARG_VGPR, ctx->ac.v2i32, &ctx->persp_center);
1180 add_arg(&args, ARG_VGPR, ctx->ac.v2i32, &ctx->persp_centroid);
1181 add_arg(&args, ARG_VGPR, ctx->ac.v3i32, NULL); /* persp pull model */
1182 add_arg(&args, ARG_VGPR, ctx->ac.v2i32, &ctx->linear_sample);
1183 add_arg(&args, ARG_VGPR, ctx->ac.v2i32, &ctx->linear_center);
1184 add_arg(&args, ARG_VGPR, ctx->ac.v2i32, &ctx->linear_centroid);
1185 add_arg(&args, ARG_VGPR, ctx->ac.f32, NULL); /* line stipple tex */
1186 add_arg(&args, ARG_VGPR, ctx->ac.f32, &ctx->abi.frag_pos[0]);
1187 add_arg(&args, ARG_VGPR, ctx->ac.f32, &ctx->abi.frag_pos[1]);
1188 add_arg(&args, ARG_VGPR, ctx->ac.f32, &ctx->abi.frag_pos[2]);
1189 add_arg(&args, ARG_VGPR, ctx->ac.f32, &ctx->abi.frag_pos[3]);
1190 add_arg(&args, ARG_VGPR, ctx->ac.i32, &ctx->abi.front_face);
1191 add_arg(&args, ARG_VGPR, ctx->ac.i32, &ctx->abi.ancillary);
1192 add_arg(&args, ARG_VGPR, ctx->ac.i32, &ctx->abi.sample_coverage);
1193 add_arg(&args, ARG_VGPR, ctx->ac.i32, NULL); /* fixed pt */
1194 break;
1195 default:
1196 unreachable("Shader stage not implemented");
1197 }
1198
1199 ctx->main_function = create_llvm_function(
1200 ctx->context, ctx->ac.module, ctx->ac.builder, NULL, 0, &args,
1201 ctx->max_workgroup_size, ctx->options);
1202 set_llvm_calling_convention(ctx->main_function, stage);
1203
1204
1205 ctx->shader_info->num_input_vgprs = 0;
1206 ctx->shader_info->num_input_sgprs = ctx->options->supports_spill ? 2 : 0;
1207
1208 ctx->shader_info->num_input_sgprs += args.num_sgprs_used;
1209
1210 if (ctx->stage != MESA_SHADER_FRAGMENT)
1211 ctx->shader_info->num_input_vgprs = args.num_vgprs_used;
1212
1213 assign_arguments(ctx->main_function, &args);
1214
1215 user_sgpr_idx = 0;
1216
1217 if (ctx->options->supports_spill || user_sgpr_info.need_ring_offsets) {
1218 set_loc_shader_ptr(ctx, AC_UD_SCRATCH_RING_OFFSETS,
1219 &user_sgpr_idx);
1220 if (ctx->options->supports_spill) {
1221 ctx->ring_offsets = ac_build_intrinsic(&ctx->ac, "llvm.amdgcn.implicit.buffer.ptr",
1222 LLVMPointerType(ctx->ac.i8, AC_ADDR_SPACE_CONST),
1223 NULL, 0, AC_FUNC_ATTR_READNONE);
1224 ctx->ring_offsets = LLVMBuildBitCast(ctx->ac.builder, ctx->ring_offsets,
1225 ac_array_in_const_addr_space(ctx->ac.v4i32), "");
1226 }
1227 }
1228
1229 /* For merged shaders the user SGPRs start at 8, with 8 system SGPRs in front (including
1230 * the rw_buffers at s0/s1. With user SGPR0 = s8, lets restart the count from 0 */
1231 if (has_previous_stage)
1232 user_sgpr_idx = 0;
1233
1234 set_global_input_locs(ctx, &user_sgpr_info, desc_sets, &user_sgpr_idx);
1235
1236 switch (stage) {
1237 case MESA_SHADER_COMPUTE:
1238 if (ctx->shader_info->info.cs.uses_grid_size) {
1239 set_loc_shader(ctx, AC_UD_CS_GRID_SIZE,
1240 &user_sgpr_idx, 3);
1241 }
1242 break;
1243 case MESA_SHADER_VERTEX:
1244 set_vs_specific_input_locs(ctx, stage, has_previous_stage,
1245 previous_stage, &user_sgpr_idx);
1246 if (ctx->abi.view_index)
1247 set_loc_shader(ctx, AC_UD_VIEW_INDEX, &user_sgpr_idx, 1);
1248 break;
1249 case MESA_SHADER_TESS_CTRL:
1250 set_vs_specific_input_locs(ctx, stage, has_previous_stage,
1251 previous_stage, &user_sgpr_idx);
1252 if (ctx->abi.view_index)
1253 set_loc_shader(ctx, AC_UD_VIEW_INDEX, &user_sgpr_idx, 1);
1254 break;
1255 case MESA_SHADER_TESS_EVAL:
1256 if (ctx->abi.view_index)
1257 set_loc_shader(ctx, AC_UD_VIEW_INDEX, &user_sgpr_idx, 1);
1258 break;
1259 case MESA_SHADER_GEOMETRY:
1260 if (has_previous_stage) {
1261 if (previous_stage == MESA_SHADER_VERTEX)
1262 set_vs_specific_input_locs(ctx, stage,
1263 has_previous_stage,
1264 previous_stage,
1265 &user_sgpr_idx);
1266 }
1267 if (ctx->abi.view_index)
1268 set_loc_shader(ctx, AC_UD_VIEW_INDEX, &user_sgpr_idx, 1);
1269 break;
1270 case MESA_SHADER_FRAGMENT:
1271 break;
1272 default:
1273 unreachable("Shader stage not implemented");
1274 }
1275
1276 if (stage == MESA_SHADER_TESS_CTRL ||
1277 (stage == MESA_SHADER_VERTEX && ctx->options->key.vs.as_ls) ||
1278 /* GFX9 has the ESGS ring buffer in LDS. */
1279 (stage == MESA_SHADER_GEOMETRY && has_previous_stage)) {
1280 ac_declare_lds_as_pointer(&ctx->ac);
1281 }
1282
1283 ctx->shader_info->num_user_sgprs = user_sgpr_idx;
1284 }
1285
1286
1287 static LLVMValueRef
1288 radv_load_resource(struct ac_shader_abi *abi, LLVMValueRef index,
1289 unsigned desc_set, unsigned binding)
1290 {
1291 struct radv_shader_context *ctx = radv_shader_context_from_abi(abi);
1292 LLVMValueRef desc_ptr = ctx->descriptor_sets[desc_set];
1293 struct radv_pipeline_layout *pipeline_layout = ctx->options->layout;
1294 struct radv_descriptor_set_layout *layout = pipeline_layout->set[desc_set].layout;
1295 unsigned base_offset = layout->binding[binding].offset;
1296 LLVMValueRef offset, stride;
1297
1298 if (layout->binding[binding].type == VK_DESCRIPTOR_TYPE_UNIFORM_BUFFER_DYNAMIC ||
1299 layout->binding[binding].type == VK_DESCRIPTOR_TYPE_STORAGE_BUFFER_DYNAMIC) {
1300 unsigned idx = pipeline_layout->set[desc_set].dynamic_offset_start +
1301 layout->binding[binding].dynamic_offset_offset;
1302 desc_ptr = ctx->abi.push_constants;
1303 base_offset = pipeline_layout->push_constant_size + 16 * idx;
1304 stride = LLVMConstInt(ctx->ac.i32, 16, false);
1305 } else
1306 stride = LLVMConstInt(ctx->ac.i32, layout->binding[binding].size, false);
1307
1308 offset = ac_build_imad(&ctx->ac, index, stride,
1309 LLVMConstInt(ctx->ac.i32, base_offset, false));
1310
1311 desc_ptr = ac_build_gep0(&ctx->ac, desc_ptr, offset);
1312 desc_ptr = ac_cast_ptr(&ctx->ac, desc_ptr, ctx->ac.v4i32);
1313 LLVMSetMetadata(desc_ptr, ctx->ac.uniform_md_kind, ctx->ac.empty_md);
1314
1315 return desc_ptr;
1316 }
1317
1318
1319 /* The offchip buffer layout for TCS->TES is
1320 *
1321 * - attribute 0 of patch 0 vertex 0
1322 * - attribute 0 of patch 0 vertex 1
1323 * - attribute 0 of patch 0 vertex 2
1324 * ...
1325 * - attribute 0 of patch 1 vertex 0
1326 * - attribute 0 of patch 1 vertex 1
1327 * ...
1328 * - attribute 1 of patch 0 vertex 0
1329 * - attribute 1 of patch 0 vertex 1
1330 * ...
1331 * - per patch attribute 0 of patch 0
1332 * - per patch attribute 0 of patch 1
1333 * ...
1334 *
1335 * Note that every attribute has 4 components.
1336 */
1337 static LLVMValueRef get_non_vertex_index_offset(struct radv_shader_context *ctx)
1338 {
1339 uint32_t num_patches = ctx->tcs_num_patches;
1340 uint32_t num_tcs_outputs;
1341 if (ctx->stage == MESA_SHADER_TESS_CTRL)
1342 num_tcs_outputs = util_last_bit64(ctx->shader_info->info.tcs.outputs_written);
1343 else
1344 num_tcs_outputs = ctx->options->key.tes.tcs_num_outputs;
1345
1346 uint32_t output_vertex_size = num_tcs_outputs * 16;
1347 uint32_t pervertex_output_patch_size = ctx->tcs_vertices_per_patch * output_vertex_size;
1348
1349 return LLVMConstInt(ctx->ac.i32, pervertex_output_patch_size * num_patches, false);
1350 }
1351
1352 static LLVMValueRef calc_param_stride(struct radv_shader_context *ctx,
1353 LLVMValueRef vertex_index)
1354 {
1355 LLVMValueRef param_stride;
1356 if (vertex_index)
1357 param_stride = LLVMConstInt(ctx->ac.i32, ctx->tcs_vertices_per_patch * ctx->tcs_num_patches, false);
1358 else
1359 param_stride = LLVMConstInt(ctx->ac.i32, ctx->tcs_num_patches, false);
1360 return param_stride;
1361 }
1362
1363 static LLVMValueRef get_tcs_tes_buffer_address(struct radv_shader_context *ctx,
1364 LLVMValueRef vertex_index,
1365 LLVMValueRef param_index)
1366 {
1367 LLVMValueRef base_addr;
1368 LLVMValueRef param_stride, constant16;
1369 LLVMValueRef rel_patch_id = get_rel_patch_id(ctx);
1370 LLVMValueRef vertices_per_patch = LLVMConstInt(ctx->ac.i32, ctx->tcs_vertices_per_patch, false);
1371 constant16 = LLVMConstInt(ctx->ac.i32, 16, false);
1372 param_stride = calc_param_stride(ctx, vertex_index);
1373 if (vertex_index) {
1374 base_addr = ac_build_imad(&ctx->ac, rel_patch_id,
1375 vertices_per_patch, vertex_index);
1376 } else {
1377 base_addr = rel_patch_id;
1378 }
1379
1380 base_addr = LLVMBuildAdd(ctx->ac.builder, base_addr,
1381 LLVMBuildMul(ctx->ac.builder, param_index,
1382 param_stride, ""), "");
1383
1384 base_addr = LLVMBuildMul(ctx->ac.builder, base_addr, constant16, "");
1385
1386 if (!vertex_index) {
1387 LLVMValueRef patch_data_offset = get_non_vertex_index_offset(ctx);
1388
1389 base_addr = LLVMBuildAdd(ctx->ac.builder, base_addr,
1390 patch_data_offset, "");
1391 }
1392 return base_addr;
1393 }
1394
1395 static LLVMValueRef get_tcs_tes_buffer_address_params(struct radv_shader_context *ctx,
1396 unsigned param,
1397 unsigned const_index,
1398 bool is_compact,
1399 LLVMValueRef vertex_index,
1400 LLVMValueRef indir_index)
1401 {
1402 LLVMValueRef param_index;
1403
1404 if (indir_index)
1405 param_index = LLVMBuildAdd(ctx->ac.builder, LLVMConstInt(ctx->ac.i32, param, false),
1406 indir_index, "");
1407 else {
1408 if (const_index && !is_compact)
1409 param += const_index;
1410 param_index = LLVMConstInt(ctx->ac.i32, param, false);
1411 }
1412 return get_tcs_tes_buffer_address(ctx, vertex_index, param_index);
1413 }
1414
1415 static LLVMValueRef
1416 get_dw_address(struct radv_shader_context *ctx,
1417 LLVMValueRef dw_addr,
1418 unsigned param,
1419 unsigned const_index,
1420 bool compact_const_index,
1421 LLVMValueRef vertex_index,
1422 LLVMValueRef stride,
1423 LLVMValueRef indir_index)
1424
1425 {
1426
1427 if (vertex_index) {
1428 dw_addr = LLVMBuildAdd(ctx->ac.builder, dw_addr,
1429 LLVMBuildMul(ctx->ac.builder,
1430 vertex_index,
1431 stride, ""), "");
1432 }
1433
1434 if (indir_index)
1435 dw_addr = LLVMBuildAdd(ctx->ac.builder, dw_addr,
1436 LLVMBuildMul(ctx->ac.builder, indir_index,
1437 LLVMConstInt(ctx->ac.i32, 4, false), ""), "");
1438 else if (const_index && !compact_const_index)
1439 dw_addr = LLVMBuildAdd(ctx->ac.builder, dw_addr,
1440 LLVMConstInt(ctx->ac.i32, const_index * 4, false), "");
1441
1442 dw_addr = LLVMBuildAdd(ctx->ac.builder, dw_addr,
1443 LLVMConstInt(ctx->ac.i32, param * 4, false), "");
1444
1445 if (const_index && compact_const_index)
1446 dw_addr = LLVMBuildAdd(ctx->ac.builder, dw_addr,
1447 LLVMConstInt(ctx->ac.i32, const_index, false), "");
1448 return dw_addr;
1449 }
1450
1451 static LLVMValueRef
1452 load_tcs_varyings(struct ac_shader_abi *abi,
1453 LLVMTypeRef type,
1454 LLVMValueRef vertex_index,
1455 LLVMValueRef indir_index,
1456 unsigned const_index,
1457 unsigned location,
1458 unsigned driver_location,
1459 unsigned component,
1460 unsigned num_components,
1461 bool is_patch,
1462 bool is_compact,
1463 bool load_input)
1464 {
1465 struct radv_shader_context *ctx = radv_shader_context_from_abi(abi);
1466 LLVMValueRef dw_addr, stride;
1467 LLVMValueRef value[4], result;
1468 unsigned param = shader_io_get_unique_index(location);
1469
1470 if (load_input) {
1471 uint32_t input_vertex_size = (ctx->tcs_num_inputs * 16) / 4;
1472 stride = LLVMConstInt(ctx->ac.i32, input_vertex_size, false);
1473 dw_addr = get_tcs_in_current_patch_offset(ctx);
1474 } else {
1475 if (!is_patch) {
1476 stride = get_tcs_out_vertex_stride(ctx);
1477 dw_addr = get_tcs_out_current_patch_offset(ctx);
1478 } else {
1479 dw_addr = get_tcs_out_current_patch_data_offset(ctx);
1480 stride = NULL;
1481 }
1482 }
1483
1484 dw_addr = get_dw_address(ctx, dw_addr, param, const_index, is_compact, vertex_index, stride,
1485 indir_index);
1486
1487 for (unsigned i = 0; i < num_components + component; i++) {
1488 value[i] = ac_lds_load(&ctx->ac, dw_addr);
1489 dw_addr = LLVMBuildAdd(ctx->ac.builder, dw_addr,
1490 ctx->ac.i32_1, "");
1491 }
1492 result = ac_build_varying_gather_values(&ctx->ac, value, num_components, component);
1493 return result;
1494 }
1495
1496 static void
1497 store_tcs_output(struct ac_shader_abi *abi,
1498 const nir_variable *var,
1499 LLVMValueRef vertex_index,
1500 LLVMValueRef param_index,
1501 unsigned const_index,
1502 LLVMValueRef src,
1503 unsigned writemask)
1504 {
1505 struct radv_shader_context *ctx = radv_shader_context_from_abi(abi);
1506 const unsigned location = var->data.location;
1507 unsigned component = var->data.location_frac;
1508 const bool is_patch = var->data.patch;
1509 const bool is_compact = var->data.compact;
1510 LLVMValueRef dw_addr;
1511 LLVMValueRef stride = NULL;
1512 LLVMValueRef buf_addr = NULL;
1513 unsigned param;
1514 bool store_lds = true;
1515
1516 if (is_patch) {
1517 if (!(ctx->tcs_patch_outputs_read & (1U << (location - VARYING_SLOT_PATCH0))))
1518 store_lds = false;
1519 } else {
1520 if (!(ctx->tcs_outputs_read & (1ULL << location)))
1521 store_lds = false;
1522 }
1523
1524 param = shader_io_get_unique_index(location);
1525 if ((location == VARYING_SLOT_CLIP_DIST0 || location == VARYING_SLOT_CLIP_DIST1) && is_compact) {
1526 const_index += component;
1527 component = 0;
1528
1529 if (const_index >= 4) {
1530 const_index -= 4;
1531 param++;
1532 }
1533 }
1534
1535 if (!is_patch) {
1536 stride = get_tcs_out_vertex_stride(ctx);
1537 dw_addr = get_tcs_out_current_patch_offset(ctx);
1538 } else {
1539 dw_addr = get_tcs_out_current_patch_data_offset(ctx);
1540 }
1541
1542 dw_addr = get_dw_address(ctx, dw_addr, param, const_index, is_compact, vertex_index, stride,
1543 param_index);
1544 buf_addr = get_tcs_tes_buffer_address_params(ctx, param, const_index, is_compact,
1545 vertex_index, param_index);
1546
1547 bool is_tess_factor = false;
1548 if (location == VARYING_SLOT_TESS_LEVEL_INNER ||
1549 location == VARYING_SLOT_TESS_LEVEL_OUTER)
1550 is_tess_factor = true;
1551
1552 unsigned base = is_compact ? const_index : 0;
1553 for (unsigned chan = 0; chan < 8; chan++) {
1554 if (!(writemask & (1 << chan)))
1555 continue;
1556 LLVMValueRef value = ac_llvm_extract_elem(&ctx->ac, src, chan - component);
1557 value = ac_to_integer(&ctx->ac, value);
1558 value = LLVMBuildZExtOrBitCast(ctx->ac.builder, value, ctx->ac.i32, "");
1559
1560 if (store_lds || is_tess_factor) {
1561 LLVMValueRef dw_addr_chan =
1562 LLVMBuildAdd(ctx->ac.builder, dw_addr,
1563 LLVMConstInt(ctx->ac.i32, chan, false), "");
1564 ac_lds_store(&ctx->ac, dw_addr_chan, value);
1565 }
1566
1567 if (!is_tess_factor && writemask != 0xF)
1568 ac_build_buffer_store_dword(&ctx->ac, ctx->hs_ring_tess_offchip, value, 1,
1569 buf_addr, ctx->oc_lds,
1570 4 * (base + chan), 1, 0, true, false);
1571 }
1572
1573 if (writemask == 0xF) {
1574 ac_build_buffer_store_dword(&ctx->ac, ctx->hs_ring_tess_offchip, src, 4,
1575 buf_addr, ctx->oc_lds,
1576 (base * 4), 1, 0, true, false);
1577 }
1578 }
1579
1580 static LLVMValueRef
1581 load_tes_input(struct ac_shader_abi *abi,
1582 LLVMTypeRef type,
1583 LLVMValueRef vertex_index,
1584 LLVMValueRef param_index,
1585 unsigned const_index,
1586 unsigned location,
1587 unsigned driver_location,
1588 unsigned component,
1589 unsigned num_components,
1590 bool is_patch,
1591 bool is_compact,
1592 bool load_input)
1593 {
1594 struct radv_shader_context *ctx = radv_shader_context_from_abi(abi);
1595 LLVMValueRef buf_addr;
1596 LLVMValueRef result;
1597 unsigned param = shader_io_get_unique_index(location);
1598
1599 if ((location == VARYING_SLOT_CLIP_DIST0 || location == VARYING_SLOT_CLIP_DIST1) && is_compact) {
1600 const_index += component;
1601 component = 0;
1602 if (const_index >= 4) {
1603 const_index -= 4;
1604 param++;
1605 }
1606 }
1607
1608 buf_addr = get_tcs_tes_buffer_address_params(ctx, param, const_index,
1609 is_compact, vertex_index, param_index);
1610
1611 LLVMValueRef comp_offset = LLVMConstInt(ctx->ac.i32, component * 4, false);
1612 buf_addr = LLVMBuildAdd(ctx->ac.builder, buf_addr, comp_offset, "");
1613
1614 result = ac_build_buffer_load(&ctx->ac, ctx->hs_ring_tess_offchip, num_components, NULL,
1615 buf_addr, ctx->oc_lds, is_compact ? (4 * const_index) : 0, 1, 0, true, false);
1616 result = ac_trim_vector(&ctx->ac, result, num_components);
1617 return result;
1618 }
1619
1620 static LLVMValueRef
1621 load_gs_input(struct ac_shader_abi *abi,
1622 unsigned location,
1623 unsigned driver_location,
1624 unsigned component,
1625 unsigned num_components,
1626 unsigned vertex_index,
1627 unsigned const_index,
1628 LLVMTypeRef type)
1629 {
1630 struct radv_shader_context *ctx = radv_shader_context_from_abi(abi);
1631 LLVMValueRef vtx_offset;
1632 unsigned param, vtx_offset_param;
1633 LLVMValueRef value[4], result;
1634
1635 vtx_offset_param = vertex_index;
1636 assert(vtx_offset_param < 6);
1637 vtx_offset = LLVMBuildMul(ctx->ac.builder, ctx->gs_vtx_offset[vtx_offset_param],
1638 LLVMConstInt(ctx->ac.i32, 4, false), "");
1639
1640 param = shader_io_get_unique_index(location);
1641
1642 for (unsigned i = component; i < num_components + component; i++) {
1643 if (ctx->ac.chip_class >= GFX9) {
1644 LLVMValueRef dw_addr = ctx->gs_vtx_offset[vtx_offset_param];
1645 dw_addr = LLVMBuildAdd(ctx->ac.builder, dw_addr,
1646 LLVMConstInt(ctx->ac.i32, param * 4 + i + const_index, 0), "");
1647 value[i] = ac_lds_load(&ctx->ac, dw_addr);
1648 } else {
1649 LLVMValueRef soffset =
1650 LLVMConstInt(ctx->ac.i32,
1651 (param * 4 + i + const_index) * 256,
1652 false);
1653
1654 value[i] = ac_build_buffer_load(&ctx->ac,
1655 ctx->esgs_ring, 1,
1656 ctx->ac.i32_0,
1657 vtx_offset, soffset,
1658 0, 1, 0, true, false);
1659 }
1660
1661 if (ac_get_type_size(type) == 2) {
1662 value[i] = LLVMBuildBitCast(ctx->ac.builder, value[i], ctx->ac.i32, "");
1663 value[i] = LLVMBuildTrunc(ctx->ac.builder, value[i], ctx->ac.i16, "");
1664 }
1665 value[i] = LLVMBuildBitCast(ctx->ac.builder, value[i], type, "");
1666 }
1667 result = ac_build_varying_gather_values(&ctx->ac, value, num_components, component);
1668 result = ac_to_integer(&ctx->ac, result);
1669 return result;
1670 }
1671
1672
1673 static void radv_emit_kill(struct ac_shader_abi *abi, LLVMValueRef visible)
1674 {
1675 struct radv_shader_context *ctx = radv_shader_context_from_abi(abi);
1676 ac_build_kill_if_false(&ctx->ac, visible);
1677 }
1678
1679 static LLVMValueRef lookup_interp_param(struct ac_shader_abi *abi,
1680 enum glsl_interp_mode interp, unsigned location)
1681 {
1682 struct radv_shader_context *ctx = radv_shader_context_from_abi(abi);
1683
1684 switch (interp) {
1685 case INTERP_MODE_FLAT:
1686 default:
1687 return NULL;
1688 case INTERP_MODE_SMOOTH:
1689 case INTERP_MODE_NONE:
1690 if (location == INTERP_CENTER)
1691 return ctx->persp_center;
1692 else if (location == INTERP_CENTROID)
1693 return ctx->persp_centroid;
1694 else if (location == INTERP_SAMPLE)
1695 return ctx->persp_sample;
1696 break;
1697 case INTERP_MODE_NOPERSPECTIVE:
1698 if (location == INTERP_CENTER)
1699 return ctx->linear_center;
1700 else if (location == INTERP_CENTROID)
1701 return ctx->linear_centroid;
1702 else if (location == INTERP_SAMPLE)
1703 return ctx->linear_sample;
1704 break;
1705 }
1706 return NULL;
1707 }
1708
1709 static uint32_t
1710 radv_get_sample_pos_offset(uint32_t num_samples)
1711 {
1712 uint32_t sample_pos_offset = 0;
1713
1714 switch (num_samples) {
1715 case 2:
1716 sample_pos_offset = 1;
1717 break;
1718 case 4:
1719 sample_pos_offset = 3;
1720 break;
1721 case 8:
1722 sample_pos_offset = 7;
1723 break;
1724 default:
1725 break;
1726 }
1727 return sample_pos_offset;
1728 }
1729
1730 static LLVMValueRef load_sample_position(struct ac_shader_abi *abi,
1731 LLVMValueRef sample_id)
1732 {
1733 struct radv_shader_context *ctx = radv_shader_context_from_abi(abi);
1734
1735 LLVMValueRef result;
1736 LLVMValueRef ptr = ac_build_gep0(&ctx->ac, ctx->ring_offsets, LLVMConstInt(ctx->ac.i32, RING_PS_SAMPLE_POSITIONS, false));
1737
1738 ptr = LLVMBuildBitCast(ctx->ac.builder, ptr,
1739 ac_array_in_const_addr_space(ctx->ac.v2f32), "");
1740
1741 uint32_t sample_pos_offset =
1742 radv_get_sample_pos_offset(ctx->options->key.fs.num_samples);
1743
1744 sample_id =
1745 LLVMBuildAdd(ctx->ac.builder, sample_id,
1746 LLVMConstInt(ctx->ac.i32, sample_pos_offset, false), "");
1747 result = ac_build_load_invariant(&ctx->ac, ptr, sample_id);
1748
1749 return result;
1750 }
1751
1752
1753 static LLVMValueRef load_sample_mask_in(struct ac_shader_abi *abi)
1754 {
1755 struct radv_shader_context *ctx = radv_shader_context_from_abi(abi);
1756 uint8_t log2_ps_iter_samples;
1757
1758 if (ctx->shader_info->info.ps.force_persample) {
1759 log2_ps_iter_samples =
1760 util_logbase2(ctx->options->key.fs.num_samples);
1761 } else {
1762 log2_ps_iter_samples = ctx->options->key.fs.log2_ps_iter_samples;
1763 }
1764
1765 /* The bit pattern matches that used by fixed function fragment
1766 * processing. */
1767 static const uint16_t ps_iter_masks[] = {
1768 0xffff, /* not used */
1769 0x5555,
1770 0x1111,
1771 0x0101,
1772 0x0001,
1773 };
1774 assert(log2_ps_iter_samples < ARRAY_SIZE(ps_iter_masks));
1775
1776 uint32_t ps_iter_mask = ps_iter_masks[log2_ps_iter_samples];
1777
1778 LLVMValueRef result, sample_id;
1779 sample_id = ac_unpack_param(&ctx->ac, abi->ancillary, 8, 4);
1780 sample_id = LLVMBuildShl(ctx->ac.builder, LLVMConstInt(ctx->ac.i32, ps_iter_mask, false), sample_id, "");
1781 result = LLVMBuildAnd(ctx->ac.builder, sample_id, abi->sample_coverage, "");
1782 return result;
1783 }
1784
1785
1786 static void
1787 visit_emit_vertex(struct ac_shader_abi *abi, unsigned stream, LLVMValueRef *addrs)
1788 {
1789 LLVMValueRef gs_next_vertex;
1790 LLVMValueRef can_emit;
1791 unsigned offset = 0;
1792 struct radv_shader_context *ctx = radv_shader_context_from_abi(abi);
1793
1794 /* Write vertex attribute values to GSVS ring */
1795 gs_next_vertex = LLVMBuildLoad(ctx->ac.builder,
1796 ctx->gs_next_vertex[stream],
1797 "");
1798
1799 /* If this thread has already emitted the declared maximum number of
1800 * vertices, kill it: excessive vertex emissions are not supposed to
1801 * have any effect, and GS threads have no externally observable
1802 * effects other than emitting vertices.
1803 */
1804 can_emit = LLVMBuildICmp(ctx->ac.builder, LLVMIntULT, gs_next_vertex,
1805 LLVMConstInt(ctx->ac.i32, ctx->gs_max_out_vertices, false), "");
1806 ac_build_kill_if_false(&ctx->ac, can_emit);
1807
1808 for (unsigned i = 0; i < AC_LLVM_MAX_OUTPUTS; ++i) {
1809 unsigned output_usage_mask =
1810 ctx->shader_info->info.gs.output_usage_mask[i];
1811 uint8_t output_stream =
1812 ctx->shader_info->info.gs.output_streams[i];
1813 LLVMValueRef *out_ptr = &addrs[i * 4];
1814 int length = util_last_bit(output_usage_mask);
1815
1816 if (!(ctx->output_mask & (1ull << i)) ||
1817 output_stream != stream)
1818 continue;
1819
1820 for (unsigned j = 0; j < length; j++) {
1821 if (!(output_usage_mask & (1 << j)))
1822 continue;
1823
1824 LLVMValueRef out_val = LLVMBuildLoad(ctx->ac.builder,
1825 out_ptr[j], "");
1826 LLVMValueRef voffset =
1827 LLVMConstInt(ctx->ac.i32, offset *
1828 ctx->gs_max_out_vertices, false);
1829
1830 offset++;
1831
1832 voffset = LLVMBuildAdd(ctx->ac.builder, voffset, gs_next_vertex, "");
1833 voffset = LLVMBuildMul(ctx->ac.builder, voffset, LLVMConstInt(ctx->ac.i32, 4, false), "");
1834
1835 out_val = ac_to_integer(&ctx->ac, out_val);
1836 out_val = LLVMBuildZExtOrBitCast(ctx->ac.builder, out_val, ctx->ac.i32, "");
1837
1838 ac_build_buffer_store_dword(&ctx->ac,
1839 ctx->gsvs_ring[stream],
1840 out_val, 1,
1841 voffset, ctx->gs2vs_offset, 0,
1842 1, 1, true, true);
1843 }
1844 }
1845
1846 gs_next_vertex = LLVMBuildAdd(ctx->ac.builder, gs_next_vertex,
1847 ctx->ac.i32_1, "");
1848 LLVMBuildStore(ctx->ac.builder, gs_next_vertex, ctx->gs_next_vertex[stream]);
1849
1850 ac_build_sendmsg(&ctx->ac,
1851 AC_SENDMSG_GS_OP_EMIT | AC_SENDMSG_GS | (stream << 8),
1852 ctx->gs_wave_id);
1853 }
1854
1855 static void
1856 visit_end_primitive(struct ac_shader_abi *abi, unsigned stream)
1857 {
1858 struct radv_shader_context *ctx = radv_shader_context_from_abi(abi);
1859 ac_build_sendmsg(&ctx->ac, AC_SENDMSG_GS_OP_CUT | AC_SENDMSG_GS | (stream << 8), ctx->gs_wave_id);
1860 }
1861
1862 static LLVMValueRef
1863 load_tess_coord(struct ac_shader_abi *abi)
1864 {
1865 struct radv_shader_context *ctx = radv_shader_context_from_abi(abi);
1866
1867 LLVMValueRef coord[4] = {
1868 ctx->tes_u,
1869 ctx->tes_v,
1870 ctx->ac.f32_0,
1871 ctx->ac.f32_0,
1872 };
1873
1874 if (ctx->tes_primitive_mode == GL_TRIANGLES)
1875 coord[2] = LLVMBuildFSub(ctx->ac.builder, ctx->ac.f32_1,
1876 LLVMBuildFAdd(ctx->ac.builder, coord[0], coord[1], ""), "");
1877
1878 return ac_build_gather_values(&ctx->ac, coord, 3);
1879 }
1880
1881 static LLVMValueRef
1882 load_patch_vertices_in(struct ac_shader_abi *abi)
1883 {
1884 struct radv_shader_context *ctx = radv_shader_context_from_abi(abi);
1885 return LLVMConstInt(ctx->ac.i32, ctx->options->key.tcs.input_vertices, false);
1886 }
1887
1888
1889 static LLVMValueRef radv_load_base_vertex(struct ac_shader_abi *abi)
1890 {
1891 return abi->base_vertex;
1892 }
1893
1894 static LLVMValueRef radv_load_ssbo(struct ac_shader_abi *abi,
1895 LLVMValueRef buffer_ptr, bool write)
1896 {
1897 struct radv_shader_context *ctx = radv_shader_context_from_abi(abi);
1898 LLVMValueRef result;
1899
1900 LLVMSetMetadata(buffer_ptr, ctx->ac.uniform_md_kind, ctx->ac.empty_md);
1901
1902 result = LLVMBuildLoad(ctx->ac.builder, buffer_ptr, "");
1903 LLVMSetMetadata(result, ctx->ac.invariant_load_md_kind, ctx->ac.empty_md);
1904
1905 return result;
1906 }
1907
1908 static LLVMValueRef radv_load_ubo(struct ac_shader_abi *abi, LLVMValueRef buffer_ptr)
1909 {
1910 struct radv_shader_context *ctx = radv_shader_context_from_abi(abi);
1911 LLVMValueRef result;
1912
1913 LLVMSetMetadata(buffer_ptr, ctx->ac.uniform_md_kind, ctx->ac.empty_md);
1914
1915 result = LLVMBuildLoad(ctx->ac.builder, buffer_ptr, "");
1916 LLVMSetMetadata(result, ctx->ac.invariant_load_md_kind, ctx->ac.empty_md);
1917
1918 return result;
1919 }
1920
1921 static LLVMValueRef radv_get_sampler_desc(struct ac_shader_abi *abi,
1922 unsigned descriptor_set,
1923 unsigned base_index,
1924 unsigned constant_index,
1925 LLVMValueRef index,
1926 enum ac_descriptor_type desc_type,
1927 bool image, bool write,
1928 bool bindless)
1929 {
1930 struct radv_shader_context *ctx = radv_shader_context_from_abi(abi);
1931 LLVMValueRef list = ctx->descriptor_sets[descriptor_set];
1932 struct radv_descriptor_set_layout *layout = ctx->options->layout->set[descriptor_set].layout;
1933 struct radv_descriptor_set_binding_layout *binding = layout->binding + base_index;
1934 unsigned offset = binding->offset;
1935 unsigned stride = binding->size;
1936 unsigned type_size;
1937 LLVMBuilderRef builder = ctx->ac.builder;
1938 LLVMTypeRef type;
1939
1940 assert(base_index < layout->binding_count);
1941
1942 switch (desc_type) {
1943 case AC_DESC_IMAGE:
1944 type = ctx->ac.v8i32;
1945 type_size = 32;
1946 break;
1947 case AC_DESC_FMASK:
1948 type = ctx->ac.v8i32;
1949 offset += 32;
1950 type_size = 32;
1951 break;
1952 case AC_DESC_SAMPLER:
1953 type = ctx->ac.v4i32;
1954 if (binding->type == VK_DESCRIPTOR_TYPE_COMBINED_IMAGE_SAMPLER)
1955 offset += 64;
1956
1957 type_size = 16;
1958 break;
1959 case AC_DESC_BUFFER:
1960 type = ctx->ac.v4i32;
1961 type_size = 16;
1962 break;
1963 default:
1964 unreachable("invalid desc_type\n");
1965 }
1966
1967 offset += constant_index * stride;
1968
1969 if (desc_type == AC_DESC_SAMPLER && binding->immutable_samplers_offset &&
1970 (!index || binding->immutable_samplers_equal)) {
1971 if (binding->immutable_samplers_equal)
1972 constant_index = 0;
1973
1974 const uint32_t *samplers = radv_immutable_samplers(layout, binding);
1975
1976 LLVMValueRef constants[] = {
1977 LLVMConstInt(ctx->ac.i32, samplers[constant_index * 4 + 0], 0),
1978 LLVMConstInt(ctx->ac.i32, samplers[constant_index * 4 + 1], 0),
1979 LLVMConstInt(ctx->ac.i32, samplers[constant_index * 4 + 2], 0),
1980 LLVMConstInt(ctx->ac.i32, samplers[constant_index * 4 + 3], 0),
1981 };
1982 return ac_build_gather_values(&ctx->ac, constants, 4);
1983 }
1984
1985 assert(stride % type_size == 0);
1986
1987 if (!index)
1988 index = ctx->ac.i32_0;
1989
1990 index = LLVMBuildMul(builder, index, LLVMConstInt(ctx->ac.i32, stride / type_size, 0), "");
1991
1992 list = ac_build_gep0(&ctx->ac, list, LLVMConstInt(ctx->ac.i32, offset, 0));
1993 list = LLVMBuildPointerCast(builder, list,
1994 ac_array_in_const32_addr_space(type), "");
1995
1996 return ac_build_load_to_sgpr(&ctx->ac, list, index);
1997 }
1998
1999 /* For 2_10_10_10 formats the alpha is handled as unsigned by pre-vega HW.
2000 * so we may need to fix it up. */
2001 static LLVMValueRef
2002 adjust_vertex_fetch_alpha(struct radv_shader_context *ctx,
2003 unsigned adjustment,
2004 LLVMValueRef alpha)
2005 {
2006 if (adjustment == RADV_ALPHA_ADJUST_NONE)
2007 return alpha;
2008
2009 LLVMValueRef c30 = LLVMConstInt(ctx->ac.i32, 30, 0);
2010
2011 alpha = LLVMBuildBitCast(ctx->ac.builder, alpha, ctx->ac.f32, "");
2012
2013 if (adjustment == RADV_ALPHA_ADJUST_SSCALED)
2014 alpha = LLVMBuildFPToUI(ctx->ac.builder, alpha, ctx->ac.i32, "");
2015 else
2016 alpha = ac_to_integer(&ctx->ac, alpha);
2017
2018 /* For the integer-like cases, do a natural sign extension.
2019 *
2020 * For the SNORM case, the values are 0.0, 0.333, 0.666, 1.0
2021 * and happen to contain 0, 1, 2, 3 as the two LSBs of the
2022 * exponent.
2023 */
2024 alpha = LLVMBuildShl(ctx->ac.builder, alpha,
2025 adjustment == RADV_ALPHA_ADJUST_SNORM ?
2026 LLVMConstInt(ctx->ac.i32, 7, 0) : c30, "");
2027 alpha = LLVMBuildAShr(ctx->ac.builder, alpha, c30, "");
2028
2029 /* Convert back to the right type. */
2030 if (adjustment == RADV_ALPHA_ADJUST_SNORM) {
2031 LLVMValueRef clamp;
2032 LLVMValueRef neg_one = LLVMConstReal(ctx->ac.f32, -1.0);
2033 alpha = LLVMBuildSIToFP(ctx->ac.builder, alpha, ctx->ac.f32, "");
2034 clamp = LLVMBuildFCmp(ctx->ac.builder, LLVMRealULT, alpha, neg_one, "");
2035 alpha = LLVMBuildSelect(ctx->ac.builder, clamp, neg_one, alpha, "");
2036 } else if (adjustment == RADV_ALPHA_ADJUST_SSCALED) {
2037 alpha = LLVMBuildSIToFP(ctx->ac.builder, alpha, ctx->ac.f32, "");
2038 }
2039
2040 return LLVMBuildBitCast(ctx->ac.builder, alpha, ctx->ac.i32, "");
2041 }
2042
2043 static unsigned
2044 get_num_channels_from_data_format(unsigned data_format)
2045 {
2046 switch (data_format) {
2047 case V_008F0C_BUF_DATA_FORMAT_8:
2048 case V_008F0C_BUF_DATA_FORMAT_16:
2049 case V_008F0C_BUF_DATA_FORMAT_32:
2050 return 1;
2051 case V_008F0C_BUF_DATA_FORMAT_8_8:
2052 case V_008F0C_BUF_DATA_FORMAT_16_16:
2053 case V_008F0C_BUF_DATA_FORMAT_32_32:
2054 return 2;
2055 case V_008F0C_BUF_DATA_FORMAT_10_11_11:
2056 case V_008F0C_BUF_DATA_FORMAT_11_11_10:
2057 case V_008F0C_BUF_DATA_FORMAT_32_32_32:
2058 return 3;
2059 case V_008F0C_BUF_DATA_FORMAT_8_8_8_8:
2060 case V_008F0C_BUF_DATA_FORMAT_10_10_10_2:
2061 case V_008F0C_BUF_DATA_FORMAT_2_10_10_10:
2062 case V_008F0C_BUF_DATA_FORMAT_16_16_16_16:
2063 case V_008F0C_BUF_DATA_FORMAT_32_32_32_32:
2064 return 4;
2065 default:
2066 break;
2067 }
2068
2069 return 4;
2070 }
2071
2072 static LLVMValueRef
2073 radv_fixup_vertex_input_fetches(struct radv_shader_context *ctx,
2074 LLVMValueRef value,
2075 unsigned num_channels,
2076 bool is_float)
2077 {
2078 LLVMValueRef zero = is_float ? ctx->ac.f32_0 : ctx->ac.i32_0;
2079 LLVMValueRef one = is_float ? ctx->ac.f32_1 : ctx->ac.i32_1;
2080 LLVMValueRef chan[4];
2081
2082 if (LLVMGetTypeKind(LLVMTypeOf(value)) == LLVMVectorTypeKind) {
2083 unsigned vec_size = LLVMGetVectorSize(LLVMTypeOf(value));
2084
2085 if (num_channels == 4 && num_channels == vec_size)
2086 return value;
2087
2088 num_channels = MIN2(num_channels, vec_size);
2089
2090 for (unsigned i = 0; i < num_channels; i++)
2091 chan[i] = ac_llvm_extract_elem(&ctx->ac, value, i);
2092 } else {
2093 if (num_channels) {
2094 assert(num_channels == 1);
2095 chan[0] = value;
2096 }
2097 }
2098
2099 for (unsigned i = num_channels; i < 4; i++) {
2100 chan[i] = i == 3 ? one : zero;
2101 chan[i] = ac_to_integer(&ctx->ac, chan[i]);
2102 }
2103
2104 return ac_build_gather_values(&ctx->ac, chan, 4);
2105 }
2106
2107 static void
2108 handle_vs_input_decl(struct radv_shader_context *ctx,
2109 struct nir_variable *variable)
2110 {
2111 LLVMValueRef t_list_ptr = ctx->vertex_buffers;
2112 LLVMValueRef t_offset;
2113 LLVMValueRef t_list;
2114 LLVMValueRef input;
2115 LLVMValueRef buffer_index;
2116 unsigned attrib_count = glsl_count_attribute_slots(variable->type, true);
2117 uint8_t input_usage_mask =
2118 ctx->shader_info->info.vs.input_usage_mask[variable->data.location];
2119 unsigned num_input_channels = util_last_bit(input_usage_mask);
2120
2121 variable->data.driver_location = variable->data.location * 4;
2122
2123 enum glsl_base_type type = glsl_get_base_type(variable->type);
2124 for (unsigned i = 0; i < attrib_count; ++i) {
2125 LLVMValueRef output[4];
2126 unsigned attrib_index = variable->data.location + i - VERT_ATTRIB_GENERIC0;
2127 unsigned attrib_format = ctx->options->key.vs.vertex_attribute_formats[attrib_index];
2128 unsigned data_format = attrib_format & 0x0f;
2129 unsigned num_format = (attrib_format >> 4) & 0x07;
2130 bool is_float = num_format != V_008F0C_BUF_NUM_FORMAT_UINT &&
2131 num_format != V_008F0C_BUF_NUM_FORMAT_SINT;
2132
2133 if (ctx->options->key.vs.instance_rate_inputs & (1u << attrib_index)) {
2134 uint32_t divisor = ctx->options->key.vs.instance_rate_divisors[attrib_index];
2135
2136 if (divisor) {
2137 buffer_index = ctx->abi.instance_id;
2138
2139 if (divisor != 1) {
2140 buffer_index = LLVMBuildUDiv(ctx->ac.builder, buffer_index,
2141 LLVMConstInt(ctx->ac.i32, divisor, 0), "");
2142 }
2143
2144 if (ctx->options->key.vs.as_ls) {
2145 ctx->shader_info->vs.vgpr_comp_cnt =
2146 MAX2(2, ctx->shader_info->vs.vgpr_comp_cnt);
2147 } else {
2148 ctx->shader_info->vs.vgpr_comp_cnt =
2149 MAX2(1, ctx->shader_info->vs.vgpr_comp_cnt);
2150 }
2151 } else {
2152 buffer_index = ctx->ac.i32_0;
2153 }
2154
2155 buffer_index = LLVMBuildAdd(ctx->ac.builder, ctx->abi.start_instance, buffer_index, "");
2156 } else
2157 buffer_index = LLVMBuildAdd(ctx->ac.builder, ctx->abi.vertex_id,
2158 ctx->abi.base_vertex, "");
2159
2160 /* Adjust the number of channels to load based on the vertex
2161 * attribute format.
2162 */
2163 unsigned num_format_channels = get_num_channels_from_data_format(data_format);
2164 unsigned num_channels = MIN2(num_input_channels, num_format_channels);
2165 unsigned attrib_binding = ctx->options->key.vs.vertex_attribute_bindings[attrib_index];
2166 unsigned attrib_offset = ctx->options->key.vs.vertex_attribute_offsets[attrib_index];
2167 unsigned attrib_stride = ctx->options->key.vs.vertex_attribute_strides[attrib_index];
2168
2169 if (ctx->options->key.vs.post_shuffle & (1 << attrib_index)) {
2170 /* Always load, at least, 3 channels for formats that
2171 * need to be shuffled because X<->Z.
2172 */
2173 num_channels = MAX2(num_channels, 3);
2174 }
2175
2176 if (attrib_stride != 0 && attrib_offset > attrib_stride) {
2177 LLVMValueRef buffer_offset =
2178 LLVMConstInt(ctx->ac.i32,
2179 attrib_offset / attrib_stride, false);
2180
2181 buffer_index = LLVMBuildAdd(ctx->ac.builder,
2182 buffer_index,
2183 buffer_offset, "");
2184
2185 attrib_offset = attrib_offset % attrib_stride;
2186 }
2187
2188 t_offset = LLVMConstInt(ctx->ac.i32, attrib_binding, false);
2189 t_list = ac_build_load_to_sgpr(&ctx->ac, t_list_ptr, t_offset);
2190
2191 input = ac_build_struct_tbuffer_load(&ctx->ac, t_list,
2192 buffer_index,
2193 LLVMConstInt(ctx->ac.i32, attrib_offset, false),
2194 ctx->ac.i32_0, ctx->ac.i32_0,
2195 num_channels,
2196 data_format, num_format,
2197 false, false, true);
2198
2199 if (ctx->options->key.vs.post_shuffle & (1 << attrib_index)) {
2200 LLVMValueRef c[4];
2201 c[0] = ac_llvm_extract_elem(&ctx->ac, input, 2);
2202 c[1] = ac_llvm_extract_elem(&ctx->ac, input, 1);
2203 c[2] = ac_llvm_extract_elem(&ctx->ac, input, 0);
2204 c[3] = ac_llvm_extract_elem(&ctx->ac, input, 3);
2205
2206 input = ac_build_gather_values(&ctx->ac, c, 4);
2207 }
2208
2209 input = radv_fixup_vertex_input_fetches(ctx, input, num_channels,
2210 is_float);
2211
2212 for (unsigned chan = 0; chan < 4; chan++) {
2213 LLVMValueRef llvm_chan = LLVMConstInt(ctx->ac.i32, chan, false);
2214 output[chan] = LLVMBuildExtractElement(ctx->ac.builder, input, llvm_chan, "");
2215 if (type == GLSL_TYPE_FLOAT16) {
2216 output[chan] = LLVMBuildBitCast(ctx->ac.builder, output[chan], ctx->ac.f32, "");
2217 output[chan] = LLVMBuildFPTrunc(ctx->ac.builder, output[chan], ctx->ac.f16, "");
2218 }
2219 }
2220
2221 unsigned alpha_adjust = (ctx->options->key.vs.alpha_adjust >> (attrib_index * 2)) & 3;
2222 output[3] = adjust_vertex_fetch_alpha(ctx, alpha_adjust, output[3]);
2223
2224 for (unsigned chan = 0; chan < 4; chan++) {
2225 output[chan] = ac_to_integer(&ctx->ac, output[chan]);
2226 if (type == GLSL_TYPE_UINT16 || type == GLSL_TYPE_INT16)
2227 output[chan] = LLVMBuildTrunc(ctx->ac.builder, output[chan], ctx->ac.i16, "");
2228
2229 ctx->inputs[ac_llvm_reg_index_soa(variable->data.location + i, chan)] = output[chan];
2230 }
2231 }
2232 }
2233
2234 static void interp_fs_input(struct radv_shader_context *ctx,
2235 unsigned attr,
2236 LLVMValueRef interp_param,
2237 LLVMValueRef prim_mask,
2238 bool float16,
2239 LLVMValueRef result[4])
2240 {
2241 LLVMValueRef attr_number;
2242 unsigned chan;
2243 LLVMValueRef i, j;
2244 bool interp = !LLVMIsUndef(interp_param);
2245
2246 attr_number = LLVMConstInt(ctx->ac.i32, attr, false);
2247
2248 /* fs.constant returns the param from the middle vertex, so it's not
2249 * really useful for flat shading. It's meant to be used for custom
2250 * interpolation (but the intrinsic can't fetch from the other two
2251 * vertices).
2252 *
2253 * Luckily, it doesn't matter, because we rely on the FLAT_SHADE state
2254 * to do the right thing. The only reason we use fs.constant is that
2255 * fs.interp cannot be used on integers, because they can be equal
2256 * to NaN.
2257 */
2258 if (interp) {
2259 interp_param = LLVMBuildBitCast(ctx->ac.builder, interp_param,
2260 ctx->ac.v2f32, "");
2261
2262 i = LLVMBuildExtractElement(ctx->ac.builder, interp_param,
2263 ctx->ac.i32_0, "");
2264 j = LLVMBuildExtractElement(ctx->ac.builder, interp_param,
2265 ctx->ac.i32_1, "");
2266 }
2267
2268 for (chan = 0; chan < 4; chan++) {
2269 LLVMValueRef llvm_chan = LLVMConstInt(ctx->ac.i32, chan, false);
2270
2271 if (interp && float16) {
2272 result[chan] = ac_build_fs_interp_f16(&ctx->ac,
2273 llvm_chan,
2274 attr_number,
2275 prim_mask, i, j);
2276 } else if (interp) {
2277 result[chan] = ac_build_fs_interp(&ctx->ac,
2278 llvm_chan,
2279 attr_number,
2280 prim_mask, i, j);
2281 } else {
2282 result[chan] = ac_build_fs_interp_mov(&ctx->ac,
2283 LLVMConstInt(ctx->ac.i32, 2, false),
2284 llvm_chan,
2285 attr_number,
2286 prim_mask);
2287 result[chan] = LLVMBuildBitCast(ctx->ac.builder, result[chan], ctx->ac.i32, "");
2288 result[chan] = LLVMBuildTruncOrBitCast(ctx->ac.builder, result[chan], float16 ? ctx->ac.i16 : ctx->ac.i32, "");
2289 }
2290 }
2291 }
2292
2293 static void mark_16bit_fs_input(struct radv_shader_context *ctx,
2294 const struct glsl_type *type,
2295 int location)
2296 {
2297 if (glsl_type_is_scalar(type) || glsl_type_is_vector(type) || glsl_type_is_matrix(type)) {
2298 unsigned attrib_count = glsl_count_attribute_slots(type, false);
2299 if (glsl_type_is_16bit(type)) {
2300 ctx->float16_shaded_mask |= ((1ull << attrib_count) - 1) << location;
2301 }
2302 } else if (glsl_type_is_array(type)) {
2303 unsigned stride = glsl_count_attribute_slots(glsl_get_array_element(type), false);
2304 for (unsigned i = 0; i < glsl_get_length(type); ++i) {
2305 mark_16bit_fs_input(ctx, glsl_get_array_element(type), location + i * stride);
2306 }
2307 } else {
2308 assert(glsl_type_is_struct_or_ifc(type));
2309 for (unsigned i = 0; i < glsl_get_length(type); i++) {
2310 mark_16bit_fs_input(ctx, glsl_get_struct_field(type, i), location);
2311 location += glsl_count_attribute_slots(glsl_get_struct_field(type, i), false);
2312 }
2313 }
2314 }
2315
2316 static void
2317 handle_fs_input_decl(struct radv_shader_context *ctx,
2318 struct nir_variable *variable)
2319 {
2320 int idx = variable->data.location;
2321 unsigned attrib_count = glsl_count_attribute_slots(variable->type, false);
2322 LLVMValueRef interp = NULL;
2323 uint64_t mask;
2324
2325 variable->data.driver_location = idx * 4;
2326
2327
2328 if (variable->data.compact) {
2329 unsigned component_count = variable->data.location_frac +
2330 glsl_get_length(variable->type);
2331 attrib_count = (component_count + 3) / 4;
2332 } else
2333 mark_16bit_fs_input(ctx, variable->type, idx);
2334
2335 mask = ((1ull << attrib_count) - 1) << variable->data.location;
2336
2337 if (glsl_get_base_type(glsl_without_array(variable->type)) == GLSL_TYPE_FLOAT ||
2338 glsl_get_base_type(glsl_without_array(variable->type)) == GLSL_TYPE_FLOAT16 ||
2339 glsl_get_base_type(glsl_without_array(variable->type)) == GLSL_TYPE_STRUCT) {
2340 unsigned interp_type;
2341 if (variable->data.sample)
2342 interp_type = INTERP_SAMPLE;
2343 else if (variable->data.centroid)
2344 interp_type = INTERP_CENTROID;
2345 else
2346 interp_type = INTERP_CENTER;
2347
2348 interp = lookup_interp_param(&ctx->abi, variable->data.interpolation, interp_type);
2349 }
2350 if (interp == NULL)
2351 interp = LLVMGetUndef(ctx->ac.i32);
2352
2353 for (unsigned i = 0; i < attrib_count; ++i)
2354 ctx->inputs[ac_llvm_reg_index_soa(idx + i, 0)] = interp;
2355
2356 ctx->input_mask |= mask;
2357 }
2358
2359 static void
2360 handle_vs_inputs(struct radv_shader_context *ctx,
2361 struct nir_shader *nir) {
2362 nir_foreach_variable(variable, &nir->inputs)
2363 handle_vs_input_decl(ctx, variable);
2364 }
2365
2366 static void
2367 prepare_interp_optimize(struct radv_shader_context *ctx,
2368 struct nir_shader *nir)
2369 {
2370 bool uses_center = false;
2371 bool uses_centroid = false;
2372 nir_foreach_variable(variable, &nir->inputs) {
2373 if (glsl_get_base_type(glsl_without_array(variable->type)) != GLSL_TYPE_FLOAT ||
2374 variable->data.sample)
2375 continue;
2376
2377 if (variable->data.centroid)
2378 uses_centroid = true;
2379 else
2380 uses_center = true;
2381 }
2382
2383 if (uses_center && uses_centroid) {
2384 LLVMValueRef sel = LLVMBuildICmp(ctx->ac.builder, LLVMIntSLT, ctx->abi.prim_mask, ctx->ac.i32_0, "");
2385 ctx->persp_centroid = LLVMBuildSelect(ctx->ac.builder, sel, ctx->persp_center, ctx->persp_centroid, "");
2386 ctx->linear_centroid = LLVMBuildSelect(ctx->ac.builder, sel, ctx->linear_center, ctx->linear_centroid, "");
2387 }
2388 }
2389
2390 static void
2391 handle_fs_inputs(struct radv_shader_context *ctx,
2392 struct nir_shader *nir)
2393 {
2394 prepare_interp_optimize(ctx, nir);
2395
2396 nir_foreach_variable(variable, &nir->inputs)
2397 handle_fs_input_decl(ctx, variable);
2398
2399 unsigned index = 0;
2400
2401 if (ctx->shader_info->info.ps.uses_input_attachments ||
2402 ctx->shader_info->info.needs_multiview_view_index) {
2403 ctx->input_mask |= 1ull << VARYING_SLOT_LAYER;
2404 ctx->inputs[ac_llvm_reg_index_soa(VARYING_SLOT_LAYER, 0)] = LLVMGetUndef(ctx->ac.i32);
2405 }
2406
2407 for (unsigned i = 0; i < RADEON_LLVM_MAX_INPUTS; ++i) {
2408 LLVMValueRef interp_param;
2409 LLVMValueRef *inputs = ctx->inputs +ac_llvm_reg_index_soa(i, 0);
2410
2411 if (!(ctx->input_mask & (1ull << i)))
2412 continue;
2413
2414 if (i >= VARYING_SLOT_VAR0 || i == VARYING_SLOT_PNTC ||
2415 i == VARYING_SLOT_PRIMITIVE_ID || i == VARYING_SLOT_LAYER) {
2416 interp_param = *inputs;
2417 bool float16 = (ctx->float16_shaded_mask >> i) & 1;
2418 interp_fs_input(ctx, index, interp_param, ctx->abi.prim_mask, float16,
2419 inputs);
2420
2421 if (LLVMIsUndef(interp_param))
2422 ctx->shader_info->fs.flat_shaded_mask |= 1u << index;
2423 if (float16)
2424 ctx->shader_info->fs.float16_shaded_mask |= 1u << index;
2425 if (i >= VARYING_SLOT_VAR0)
2426 ctx->abi.fs_input_attr_indices[i - VARYING_SLOT_VAR0] = index;
2427 ++index;
2428 } else if (i == VARYING_SLOT_CLIP_DIST0) {
2429 int length = ctx->shader_info->info.ps.num_input_clips_culls;
2430
2431 for (unsigned j = 0; j < length; j += 4) {
2432 inputs = ctx->inputs + ac_llvm_reg_index_soa(i, j);
2433
2434 interp_param = *inputs;
2435 interp_fs_input(ctx, index, interp_param,
2436 ctx->abi.prim_mask, false, inputs);
2437 ++index;
2438 }
2439 } else if (i == VARYING_SLOT_POS) {
2440 for(int i = 0; i < 3; ++i)
2441 inputs[i] = ctx->abi.frag_pos[i];
2442
2443 inputs[3] = ac_build_fdiv(&ctx->ac, ctx->ac.f32_1,
2444 ctx->abi.frag_pos[3]);
2445 }
2446 }
2447 ctx->shader_info->fs.num_interp = index;
2448 ctx->shader_info->fs.input_mask = ctx->input_mask >> VARYING_SLOT_VAR0;
2449
2450 if (ctx->shader_info->info.needs_multiview_view_index)
2451 ctx->abi.view_index = ctx->inputs[ac_llvm_reg_index_soa(VARYING_SLOT_LAYER, 0)];
2452 }
2453
2454 static void
2455 scan_shader_output_decl(struct radv_shader_context *ctx,
2456 struct nir_variable *variable,
2457 struct nir_shader *shader,
2458 gl_shader_stage stage)
2459 {
2460 int idx = variable->data.location + variable->data.index;
2461 unsigned attrib_count = glsl_count_attribute_slots(variable->type, false);
2462 uint64_t mask_attribs;
2463
2464 variable->data.driver_location = idx * 4;
2465
2466 /* tess ctrl has it's own load/store paths for outputs */
2467 if (stage == MESA_SHADER_TESS_CTRL)
2468 return;
2469
2470 if (variable->data.compact) {
2471 unsigned component_count = variable->data.location_frac +
2472 glsl_get_length(variable->type);
2473 attrib_count = (component_count + 3) / 4;
2474 }
2475
2476 mask_attribs = ((1ull << attrib_count) - 1) << idx;
2477 if (stage == MESA_SHADER_VERTEX ||
2478 stage == MESA_SHADER_TESS_EVAL ||
2479 stage == MESA_SHADER_GEOMETRY) {
2480 if (idx == VARYING_SLOT_CLIP_DIST0) {
2481 if (stage == MESA_SHADER_VERTEX) {
2482 ctx->shader_info->vs.outinfo.clip_dist_mask = (1 << shader->info.clip_distance_array_size) - 1;
2483 ctx->shader_info->vs.outinfo.cull_dist_mask = (1 << shader->info.cull_distance_array_size) - 1;
2484 ctx->shader_info->vs.outinfo.cull_dist_mask <<= shader->info.clip_distance_array_size;
2485 }
2486 if (stage == MESA_SHADER_TESS_EVAL) {
2487 ctx->shader_info->tes.outinfo.clip_dist_mask = (1 << shader->info.clip_distance_array_size) - 1;
2488 ctx->shader_info->tes.outinfo.cull_dist_mask = (1 << shader->info.cull_distance_array_size) - 1;
2489 ctx->shader_info->tes.outinfo.cull_dist_mask <<= shader->info.clip_distance_array_size;
2490 }
2491 }
2492 }
2493
2494 ctx->output_mask |= mask_attribs;
2495 }
2496
2497
2498 /* Initialize arguments for the shader export intrinsic */
2499 static void
2500 si_llvm_init_export_args(struct radv_shader_context *ctx,
2501 LLVMValueRef *values,
2502 unsigned enabled_channels,
2503 unsigned target,
2504 struct ac_export_args *args)
2505 {
2506 /* Specify the channels that are enabled. */
2507 args->enabled_channels = enabled_channels;
2508
2509 /* Specify whether the EXEC mask represents the valid mask */
2510 args->valid_mask = 0;
2511
2512 /* Specify whether this is the last export */
2513 args->done = 0;
2514
2515 /* Specify the target we are exporting */
2516 args->target = target;
2517
2518 args->compr = false;
2519 args->out[0] = LLVMGetUndef(ctx->ac.f32);
2520 args->out[1] = LLVMGetUndef(ctx->ac.f32);
2521 args->out[2] = LLVMGetUndef(ctx->ac.f32);
2522 args->out[3] = LLVMGetUndef(ctx->ac.f32);
2523
2524 if (!values)
2525 return;
2526
2527 bool is_16bit = ac_get_type_size(LLVMTypeOf(values[0])) == 2;
2528 if (ctx->stage == MESA_SHADER_FRAGMENT) {
2529 unsigned index = target - V_008DFC_SQ_EXP_MRT;
2530 unsigned col_format = (ctx->options->key.fs.col_format >> (4 * index)) & 0xf;
2531 bool is_int8 = (ctx->options->key.fs.is_int8 >> index) & 1;
2532 bool is_int10 = (ctx->options->key.fs.is_int10 >> index) & 1;
2533 unsigned chan;
2534
2535 LLVMValueRef (*packf)(struct ac_llvm_context *ctx, LLVMValueRef args[2]) = NULL;
2536 LLVMValueRef (*packi)(struct ac_llvm_context *ctx, LLVMValueRef args[2],
2537 unsigned bits, bool hi) = NULL;
2538
2539 switch(col_format) {
2540 case V_028714_SPI_SHADER_ZERO:
2541 args->enabled_channels = 0; /* writemask */
2542 args->target = V_008DFC_SQ_EXP_NULL;
2543 break;
2544
2545 case V_028714_SPI_SHADER_32_R:
2546 args->enabled_channels = 1;
2547 args->out[0] = values[0];
2548 break;
2549
2550 case V_028714_SPI_SHADER_32_GR:
2551 args->enabled_channels = 0x3;
2552 args->out[0] = values[0];
2553 args->out[1] = values[1];
2554 break;
2555
2556 case V_028714_SPI_SHADER_32_AR:
2557 args->enabled_channels = 0x9;
2558 args->out[0] = values[0];
2559 args->out[3] = values[3];
2560 break;
2561
2562 case V_028714_SPI_SHADER_FP16_ABGR:
2563 args->enabled_channels = 0x5;
2564 packf = ac_build_cvt_pkrtz_f16;
2565 if (is_16bit) {
2566 for (unsigned chan = 0; chan < 4; chan++)
2567 values[chan] = LLVMBuildFPExt(ctx->ac.builder,
2568 values[chan],
2569 ctx->ac.f32, "");
2570 }
2571 break;
2572
2573 case V_028714_SPI_SHADER_UNORM16_ABGR:
2574 args->enabled_channels = 0x5;
2575 packf = ac_build_cvt_pknorm_u16;
2576 break;
2577
2578 case V_028714_SPI_SHADER_SNORM16_ABGR:
2579 args->enabled_channels = 0x5;
2580 packf = ac_build_cvt_pknorm_i16;
2581 break;
2582
2583 case V_028714_SPI_SHADER_UINT16_ABGR:
2584 args->enabled_channels = 0x5;
2585 packi = ac_build_cvt_pk_u16;
2586 if (is_16bit) {
2587 for (unsigned chan = 0; chan < 4; chan++)
2588 values[chan] = LLVMBuildZExt(ctx->ac.builder,
2589 ac_to_integer(&ctx->ac, values[chan]),
2590 ctx->ac.i32, "");
2591 }
2592 break;
2593
2594 case V_028714_SPI_SHADER_SINT16_ABGR:
2595 args->enabled_channels = 0x5;
2596 packi = ac_build_cvt_pk_i16;
2597 if (is_16bit) {
2598 for (unsigned chan = 0; chan < 4; chan++)
2599 values[chan] = LLVMBuildSExt(ctx->ac.builder,
2600 ac_to_integer(&ctx->ac, values[chan]),
2601 ctx->ac.i32, "");
2602 }
2603 break;
2604
2605 default:
2606 case V_028714_SPI_SHADER_32_ABGR:
2607 memcpy(&args->out[0], values, sizeof(values[0]) * 4);
2608 break;
2609 }
2610
2611 /* Pack f16 or norm_i16/u16. */
2612 if (packf) {
2613 for (chan = 0; chan < 2; chan++) {
2614 LLVMValueRef pack_args[2] = {
2615 values[2 * chan],
2616 values[2 * chan + 1]
2617 };
2618 LLVMValueRef packed;
2619
2620 packed = packf(&ctx->ac, pack_args);
2621 args->out[chan] = ac_to_float(&ctx->ac, packed);
2622 }
2623 args->compr = 1; /* COMPR flag */
2624 }
2625
2626 /* Pack i16/u16. */
2627 if (packi) {
2628 for (chan = 0; chan < 2; chan++) {
2629 LLVMValueRef pack_args[2] = {
2630 ac_to_integer(&ctx->ac, values[2 * chan]),
2631 ac_to_integer(&ctx->ac, values[2 * chan + 1])
2632 };
2633 LLVMValueRef packed;
2634
2635 packed = packi(&ctx->ac, pack_args,
2636 is_int8 ? 8 : is_int10 ? 10 : 16,
2637 chan == 1);
2638 args->out[chan] = ac_to_float(&ctx->ac, packed);
2639 }
2640 args->compr = 1; /* COMPR flag */
2641 }
2642 return;
2643 }
2644
2645 if (is_16bit) {
2646 for (unsigned chan = 0; chan < 4; chan++) {
2647 values[chan] = LLVMBuildBitCast(ctx->ac.builder, values[chan], ctx->ac.i16, "");
2648 args->out[chan] = LLVMBuildZExt(ctx->ac.builder, values[chan], ctx->ac.i32, "");
2649 }
2650 } else
2651 memcpy(&args->out[0], values, sizeof(values[0]) * 4);
2652
2653 for (unsigned i = 0; i < 4; ++i)
2654 args->out[i] = ac_to_float(&ctx->ac, args->out[i]);
2655 }
2656
2657 static void
2658 radv_export_param(struct radv_shader_context *ctx, unsigned index,
2659 LLVMValueRef *values, unsigned enabled_channels)
2660 {
2661 struct ac_export_args args;
2662
2663 si_llvm_init_export_args(ctx, values, enabled_channels,
2664 V_008DFC_SQ_EXP_PARAM + index, &args);
2665 ac_build_export(&ctx->ac, &args);
2666 }
2667
2668 static LLVMValueRef
2669 radv_load_output(struct radv_shader_context *ctx, unsigned index, unsigned chan)
2670 {
2671 LLVMValueRef output =
2672 ctx->abi.outputs[ac_llvm_reg_index_soa(index, chan)];
2673
2674 return LLVMBuildLoad(ctx->ac.builder, output, "");
2675 }
2676
2677 static void
2678 radv_emit_stream_output(struct radv_shader_context *ctx,
2679 LLVMValueRef const *so_buffers,
2680 LLVMValueRef const *so_write_offsets,
2681 const struct radv_stream_output *output)
2682 {
2683 unsigned num_comps = util_bitcount(output->component_mask);
2684 unsigned loc = output->location;
2685 unsigned buf = output->buffer;
2686 unsigned offset = output->offset;
2687 unsigned start;
2688 LLVMValueRef out[4];
2689
2690 assert(num_comps && num_comps <= 4);
2691 if (!num_comps || num_comps > 4)
2692 return;
2693
2694 /* Get the first component. */
2695 start = ffs(output->component_mask) - 1;
2696
2697 /* Load the output as int. */
2698 for (int i = 0; i < num_comps; i++) {
2699 out[i] = ac_to_integer(&ctx->ac,
2700 radv_load_output(ctx, loc, start + i));
2701 }
2702
2703 /* Pack the output. */
2704 LLVMValueRef vdata = NULL;
2705
2706 switch (num_comps) {
2707 case 1: /* as i32 */
2708 vdata = out[0];
2709 break;
2710 case 2: /* as v2i32 */
2711 case 3: /* as v4i32 (aligned to 4) */
2712 out[3] = LLVMGetUndef(ctx->ac.i32);
2713 /* fall through */
2714 case 4: /* as v4i32 */
2715 vdata = ac_build_gather_values(&ctx->ac, out,
2716 util_next_power_of_two(num_comps));
2717 break;
2718 }
2719
2720 ac_build_buffer_store_dword(&ctx->ac, so_buffers[buf],
2721 vdata, num_comps, so_write_offsets[buf],
2722 ctx->ac.i32_0, offset,
2723 1, 1, true, false);
2724 }
2725
2726 static void
2727 radv_emit_streamout(struct radv_shader_context *ctx, unsigned stream)
2728 {
2729 struct ac_build_if_state if_ctx;
2730 int i;
2731
2732 /* Get bits [22:16], i.e. (so_param >> 16) & 127; */
2733 assert(ctx->streamout_config);
2734 LLVMValueRef so_vtx_count =
2735 ac_build_bfe(&ctx->ac, ctx->streamout_config,
2736 LLVMConstInt(ctx->ac.i32, 16, false),
2737 LLVMConstInt(ctx->ac.i32, 7, false), false);
2738
2739 LLVMValueRef tid = ac_get_thread_id(&ctx->ac);
2740
2741 /* can_emit = tid < so_vtx_count; */
2742 LLVMValueRef can_emit = LLVMBuildICmp(ctx->ac.builder, LLVMIntULT,
2743 tid, so_vtx_count, "");
2744
2745 /* Emit the streamout code conditionally. This actually avoids
2746 * out-of-bounds buffer access. The hw tells us via the SGPR
2747 * (so_vtx_count) which threads are allowed to emit streamout data.
2748 */
2749 ac_nir_build_if(&if_ctx, ctx, can_emit);
2750 {
2751 /* The buffer offset is computed as follows:
2752 * ByteOffset = streamout_offset[buffer_id]*4 +
2753 * (streamout_write_index + thread_id)*stride[buffer_id] +
2754 * attrib_offset
2755 */
2756 LLVMValueRef so_write_index = ctx->streamout_write_idx;
2757
2758 /* Compute (streamout_write_index + thread_id). */
2759 so_write_index =
2760 LLVMBuildAdd(ctx->ac.builder, so_write_index, tid, "");
2761
2762 /* Load the descriptor and compute the write offset for each
2763 * enabled buffer.
2764 */
2765 LLVMValueRef so_write_offset[4] = {};
2766 LLVMValueRef so_buffers[4] = {};
2767 LLVMValueRef buf_ptr = ctx->streamout_buffers;
2768
2769 for (i = 0; i < 4; i++) {
2770 uint16_t stride = ctx->shader_info->info.so.strides[i];
2771
2772 if (!stride)
2773 continue;
2774
2775 LLVMValueRef offset =
2776 LLVMConstInt(ctx->ac.i32, i, false);
2777
2778 so_buffers[i] = ac_build_load_to_sgpr(&ctx->ac,
2779 buf_ptr, offset);
2780
2781 LLVMValueRef so_offset = ctx->streamout_offset[i];
2782
2783 so_offset = LLVMBuildMul(ctx->ac.builder, so_offset,
2784 LLVMConstInt(ctx->ac.i32, 4, false), "");
2785
2786 so_write_offset[i] =
2787 ac_build_imad(&ctx->ac, so_write_index,
2788 LLVMConstInt(ctx->ac.i32,
2789 stride * 4, false),
2790 so_offset);
2791 }
2792
2793 /* Write streamout data. */
2794 for (i = 0; i < ctx->shader_info->info.so.num_outputs; i++) {
2795 struct radv_stream_output *output =
2796 &ctx->shader_info->info.so.outputs[i];
2797
2798 if (stream != output->stream)
2799 continue;
2800
2801 radv_emit_stream_output(ctx, so_buffers,
2802 so_write_offset, output);
2803 }
2804 }
2805 ac_nir_build_endif(&if_ctx);
2806 }
2807
2808 static void
2809 handle_vs_outputs_post(struct radv_shader_context *ctx,
2810 bool export_prim_id, bool export_layer_id,
2811 struct radv_vs_output_info *outinfo)
2812 {
2813 uint32_t param_count = 0;
2814 unsigned target;
2815 unsigned pos_idx, num_pos_exports = 0;
2816 struct ac_export_args args, pos_args[4] = {};
2817 LLVMValueRef psize_value = NULL, layer_value = NULL, viewport_index_value = NULL;
2818 int i;
2819
2820 if (ctx->options->key.has_multiview_view_index) {
2821 LLVMValueRef* tmp_out = &ctx->abi.outputs[ac_llvm_reg_index_soa(VARYING_SLOT_LAYER, 0)];
2822 if(!*tmp_out) {
2823 for(unsigned i = 0; i < 4; ++i)
2824 ctx->abi.outputs[ac_llvm_reg_index_soa(VARYING_SLOT_LAYER, i)] =
2825 ac_build_alloca_undef(&ctx->ac, ctx->ac.f32, "");
2826 }
2827
2828 LLVMBuildStore(ctx->ac.builder, ac_to_float(&ctx->ac, ctx->abi.view_index), *tmp_out);
2829 ctx->output_mask |= 1ull << VARYING_SLOT_LAYER;
2830 }
2831
2832 memset(outinfo->vs_output_param_offset, AC_EXP_PARAM_UNDEFINED,
2833 sizeof(outinfo->vs_output_param_offset));
2834
2835 for(unsigned location = VARYING_SLOT_CLIP_DIST0; location <= VARYING_SLOT_CLIP_DIST1; ++location) {
2836 if (ctx->output_mask & (1ull << location)) {
2837 unsigned output_usage_mask, length;
2838 LLVMValueRef slots[4];
2839 unsigned j;
2840
2841 if (ctx->stage == MESA_SHADER_VERTEX &&
2842 !ctx->is_gs_copy_shader) {
2843 output_usage_mask =
2844 ctx->shader_info->info.vs.output_usage_mask[location];
2845 } else if (ctx->stage == MESA_SHADER_TESS_EVAL) {
2846 output_usage_mask =
2847 ctx->shader_info->info.tes.output_usage_mask[location];
2848 } else {
2849 assert(ctx->is_gs_copy_shader);
2850 output_usage_mask =
2851 ctx->shader_info->info.gs.output_usage_mask[location];
2852 }
2853
2854 length = util_last_bit(output_usage_mask);
2855
2856 for (j = 0; j < length; j++)
2857 slots[j] = ac_to_float(&ctx->ac, radv_load_output(ctx, location, j));
2858
2859 for (i = length; i < 4; i++)
2860 slots[i] = LLVMGetUndef(ctx->ac.f32);
2861
2862 target = V_008DFC_SQ_EXP_POS + 2 + (location - VARYING_SLOT_CLIP_DIST0);
2863 si_llvm_init_export_args(ctx, &slots[0], 0xf, target, &args);
2864 memcpy(&pos_args[target - V_008DFC_SQ_EXP_POS],
2865 &args, sizeof(args));
2866
2867 /* Export the clip/cull distances values to the next stage. */
2868 radv_export_param(ctx, param_count, &slots[0], 0xf);
2869 outinfo->vs_output_param_offset[location] = param_count++;
2870 }
2871 }
2872
2873 LLVMValueRef pos_values[4] = {ctx->ac.f32_0, ctx->ac.f32_0, ctx->ac.f32_0, ctx->ac.f32_1};
2874 if (ctx->output_mask & (1ull << VARYING_SLOT_POS)) {
2875 for (unsigned j = 0; j < 4; j++)
2876 pos_values[j] = radv_load_output(ctx, VARYING_SLOT_POS, j);
2877 }
2878 si_llvm_init_export_args(ctx, pos_values, 0xf, V_008DFC_SQ_EXP_POS, &pos_args[0]);
2879
2880 if (ctx->output_mask & (1ull << VARYING_SLOT_PSIZ)) {
2881 outinfo->writes_pointsize = true;
2882 psize_value = radv_load_output(ctx, VARYING_SLOT_PSIZ, 0);
2883 }
2884
2885 if (ctx->output_mask & (1ull << VARYING_SLOT_LAYER)) {
2886 outinfo->writes_layer = true;
2887 layer_value = radv_load_output(ctx, VARYING_SLOT_LAYER, 0);
2888 }
2889
2890 if (ctx->output_mask & (1ull << VARYING_SLOT_VIEWPORT)) {
2891 outinfo->writes_viewport_index = true;
2892 viewport_index_value = radv_load_output(ctx, VARYING_SLOT_VIEWPORT, 0);
2893 }
2894
2895 if (ctx->shader_info->info.so.num_outputs &&
2896 !ctx->is_gs_copy_shader) {
2897 /* The GS copy shader emission already emits streamout. */
2898 radv_emit_streamout(ctx, 0);
2899 }
2900
2901 if (outinfo->writes_pointsize ||
2902 outinfo->writes_layer ||
2903 outinfo->writes_viewport_index) {
2904 pos_args[1].enabled_channels = ((outinfo->writes_pointsize == true ? 1 : 0) |
2905 (outinfo->writes_layer == true ? 4 : 0));
2906 pos_args[1].valid_mask = 0;
2907 pos_args[1].done = 0;
2908 pos_args[1].target = V_008DFC_SQ_EXP_POS + 1;
2909 pos_args[1].compr = 0;
2910 pos_args[1].out[0] = ctx->ac.f32_0; /* X */
2911 pos_args[1].out[1] = ctx->ac.f32_0; /* Y */
2912 pos_args[1].out[2] = ctx->ac.f32_0; /* Z */
2913 pos_args[1].out[3] = ctx->ac.f32_0; /* W */
2914
2915 if (outinfo->writes_pointsize == true)
2916 pos_args[1].out[0] = psize_value;
2917 if (outinfo->writes_layer == true)
2918 pos_args[1].out[2] = layer_value;
2919 if (outinfo->writes_viewport_index == true) {
2920 if (ctx->options->chip_class >= GFX9) {
2921 /* GFX9 has the layer in out.z[10:0] and the viewport
2922 * index in out.z[19:16].
2923 */
2924 LLVMValueRef v = viewport_index_value;
2925 v = ac_to_integer(&ctx->ac, v);
2926 v = LLVMBuildShl(ctx->ac.builder, v,
2927 LLVMConstInt(ctx->ac.i32, 16, false),
2928 "");
2929 v = LLVMBuildOr(ctx->ac.builder, v,
2930 ac_to_integer(&ctx->ac, pos_args[1].out[2]), "");
2931
2932 pos_args[1].out[2] = ac_to_float(&ctx->ac, v);
2933 pos_args[1].enabled_channels |= 1 << 2;
2934 } else {
2935 pos_args[1].out[3] = viewport_index_value;
2936 pos_args[1].enabled_channels |= 1 << 3;
2937 }
2938 }
2939 }
2940 for (i = 0; i < 4; i++) {
2941 if (pos_args[i].out[0])
2942 num_pos_exports++;
2943 }
2944
2945 pos_idx = 0;
2946 for (i = 0; i < 4; i++) {
2947 if (!pos_args[i].out[0])
2948 continue;
2949
2950 /* Specify the target we are exporting */
2951 pos_args[i].target = V_008DFC_SQ_EXP_POS + pos_idx++;
2952 if (pos_idx == num_pos_exports)
2953 pos_args[i].done = 1;
2954 ac_build_export(&ctx->ac, &pos_args[i]);
2955 }
2956
2957 for (unsigned i = 0; i < AC_LLVM_MAX_OUTPUTS; ++i) {
2958 LLVMValueRef values[4];
2959 if (!(ctx->output_mask & (1ull << i)))
2960 continue;
2961
2962 if (i != VARYING_SLOT_LAYER &&
2963 i != VARYING_SLOT_PRIMITIVE_ID &&
2964 i < VARYING_SLOT_VAR0)
2965 continue;
2966
2967 for (unsigned j = 0; j < 4; j++)
2968 values[j] = ac_to_float(&ctx->ac, radv_load_output(ctx, i, j));
2969
2970 unsigned output_usage_mask;
2971
2972 if (ctx->stage == MESA_SHADER_VERTEX &&
2973 !ctx->is_gs_copy_shader) {
2974 output_usage_mask =
2975 ctx->shader_info->info.vs.output_usage_mask[i];
2976 } else if (ctx->stage == MESA_SHADER_TESS_EVAL) {
2977 output_usage_mask =
2978 ctx->shader_info->info.tes.output_usage_mask[i];
2979 } else {
2980 assert(ctx->is_gs_copy_shader);
2981 output_usage_mask =
2982 ctx->shader_info->info.gs.output_usage_mask[i];
2983 }
2984
2985 radv_export_param(ctx, param_count, values, output_usage_mask);
2986
2987 outinfo->vs_output_param_offset[i] = param_count++;
2988 }
2989
2990 if (export_prim_id) {
2991 LLVMValueRef values[4];
2992
2993 values[0] = ctx->vs_prim_id;
2994 ctx->shader_info->vs.vgpr_comp_cnt = MAX2(2,
2995 ctx->shader_info->vs.vgpr_comp_cnt);
2996 for (unsigned j = 1; j < 4; j++)
2997 values[j] = ctx->ac.f32_0;
2998
2999 radv_export_param(ctx, param_count, values, 0x1);
3000
3001 outinfo->vs_output_param_offset[VARYING_SLOT_PRIMITIVE_ID] = param_count++;
3002 outinfo->export_prim_id = true;
3003 }
3004
3005 if (export_layer_id && layer_value) {
3006 LLVMValueRef values[4];
3007
3008 values[0] = layer_value;
3009 for (unsigned j = 1; j < 4; j++)
3010 values[j] = ctx->ac.f32_0;
3011
3012 radv_export_param(ctx, param_count, values, 0x1);
3013
3014 outinfo->vs_output_param_offset[VARYING_SLOT_LAYER] = param_count++;
3015 }
3016
3017 outinfo->pos_exports = num_pos_exports;
3018 outinfo->param_exports = param_count;
3019 }
3020
3021 static void
3022 handle_es_outputs_post(struct radv_shader_context *ctx,
3023 struct radv_es_output_info *outinfo)
3024 {
3025 int j;
3026 uint64_t max_output_written = 0;
3027 LLVMValueRef lds_base = NULL;
3028
3029 for (unsigned i = 0; i < AC_LLVM_MAX_OUTPUTS; ++i) {
3030 int param_index;
3031
3032 if (!(ctx->output_mask & (1ull << i)))
3033 continue;
3034
3035 param_index = shader_io_get_unique_index(i);
3036
3037 max_output_written = MAX2(param_index, max_output_written);
3038 }
3039
3040 outinfo->esgs_itemsize = (max_output_written + 1) * 16;
3041
3042 if (ctx->ac.chip_class >= GFX9) {
3043 unsigned itemsize_dw = outinfo->esgs_itemsize / 4;
3044 LLVMValueRef vertex_idx = ac_get_thread_id(&ctx->ac);
3045 LLVMValueRef wave_idx = ac_unpack_param(&ctx->ac, ctx->merged_wave_info, 24, 4);
3046 vertex_idx = LLVMBuildOr(ctx->ac.builder, vertex_idx,
3047 LLVMBuildMul(ctx->ac.builder, wave_idx,
3048 LLVMConstInt(ctx->ac.i32, 64, false), ""), "");
3049 lds_base = LLVMBuildMul(ctx->ac.builder, vertex_idx,
3050 LLVMConstInt(ctx->ac.i32, itemsize_dw, 0), "");
3051 }
3052
3053 for (unsigned i = 0; i < AC_LLVM_MAX_OUTPUTS; ++i) {
3054 LLVMValueRef dw_addr = NULL;
3055 LLVMValueRef *out_ptr = &ctx->abi.outputs[i * 4];
3056 unsigned output_usage_mask;
3057 int param_index;
3058
3059 if (!(ctx->output_mask & (1ull << i)))
3060 continue;
3061
3062 if (ctx->stage == MESA_SHADER_VERTEX) {
3063 output_usage_mask =
3064 ctx->shader_info->info.vs.output_usage_mask[i];
3065 } else {
3066 assert(ctx->stage == MESA_SHADER_TESS_EVAL);
3067 output_usage_mask =
3068 ctx->shader_info->info.tes.output_usage_mask[i];
3069 }
3070
3071 param_index = shader_io_get_unique_index(i);
3072
3073 if (lds_base) {
3074 dw_addr = LLVMBuildAdd(ctx->ac.builder, lds_base,
3075 LLVMConstInt(ctx->ac.i32, param_index * 4, false),
3076 "");
3077 }
3078
3079 for (j = 0; j < 4; j++) {
3080 if (!(output_usage_mask & (1 << j)))
3081 continue;
3082
3083 LLVMValueRef out_val = LLVMBuildLoad(ctx->ac.builder, out_ptr[j], "");
3084 out_val = ac_to_integer(&ctx->ac, out_val);
3085 out_val = LLVMBuildZExtOrBitCast(ctx->ac.builder, out_val, ctx->ac.i32, "");
3086
3087 if (ctx->ac.chip_class >= GFX9) {
3088 LLVMValueRef dw_addr_offset =
3089 LLVMBuildAdd(ctx->ac.builder, dw_addr,
3090 LLVMConstInt(ctx->ac.i32,
3091 j, false), "");
3092
3093 ac_lds_store(&ctx->ac, dw_addr_offset, out_val);
3094 } else {
3095 ac_build_buffer_store_dword(&ctx->ac,
3096 ctx->esgs_ring,
3097 out_val, 1,
3098 NULL, ctx->es2gs_offset,
3099 (4 * param_index + j) * 4,
3100 1, 1, true, true);
3101 }
3102 }
3103 }
3104 }
3105
3106 static void
3107 handle_ls_outputs_post(struct radv_shader_context *ctx)
3108 {
3109 LLVMValueRef vertex_id = ctx->rel_auto_id;
3110 uint32_t num_tcs_inputs = util_last_bit64(ctx->shader_info->info.vs.ls_outputs_written);
3111 LLVMValueRef vertex_dw_stride = LLVMConstInt(ctx->ac.i32, num_tcs_inputs * 4, false);
3112 LLVMValueRef base_dw_addr = LLVMBuildMul(ctx->ac.builder, vertex_id,
3113 vertex_dw_stride, "");
3114
3115 for (unsigned i = 0; i < AC_LLVM_MAX_OUTPUTS; ++i) {
3116 LLVMValueRef *out_ptr = &ctx->abi.outputs[i * 4];
3117
3118 if (!(ctx->output_mask & (1ull << i)))
3119 continue;
3120
3121 int param = shader_io_get_unique_index(i);
3122 LLVMValueRef dw_addr = LLVMBuildAdd(ctx->ac.builder, base_dw_addr,
3123 LLVMConstInt(ctx->ac.i32, param * 4, false),
3124 "");
3125 for (unsigned j = 0; j < 4; j++) {
3126 LLVMValueRef value = LLVMBuildLoad(ctx->ac.builder, out_ptr[j], "");
3127 value = ac_to_integer(&ctx->ac, value);
3128 value = LLVMBuildZExtOrBitCast(ctx->ac.builder, value, ctx->ac.i32, "");
3129 ac_lds_store(&ctx->ac, dw_addr, value);
3130 dw_addr = LLVMBuildAdd(ctx->ac.builder, dw_addr, ctx->ac.i32_1, "");
3131 }
3132 }
3133 }
3134
3135 static void
3136 write_tess_factors(struct radv_shader_context *ctx)
3137 {
3138 unsigned stride, outer_comps, inner_comps;
3139 struct ac_build_if_state if_ctx, inner_if_ctx;
3140 LLVMValueRef invocation_id = ac_unpack_param(&ctx->ac, ctx->abi.tcs_rel_ids, 8, 5);
3141 LLVMValueRef rel_patch_id = ac_unpack_param(&ctx->ac, ctx->abi.tcs_rel_ids, 0, 8);
3142 unsigned tess_inner_index = 0, tess_outer_index;
3143 LLVMValueRef lds_base, lds_inner = NULL, lds_outer, byteoffset, buffer;
3144 LLVMValueRef out[6], vec0, vec1, tf_base, inner[4], outer[4];
3145 int i;
3146 ac_emit_barrier(&ctx->ac, ctx->stage);
3147
3148 switch (ctx->options->key.tcs.primitive_mode) {
3149 case GL_ISOLINES:
3150 stride = 2;
3151 outer_comps = 2;
3152 inner_comps = 0;
3153 break;
3154 case GL_TRIANGLES:
3155 stride = 4;
3156 outer_comps = 3;
3157 inner_comps = 1;
3158 break;
3159 case GL_QUADS:
3160 stride = 6;
3161 outer_comps = 4;
3162 inner_comps = 2;
3163 break;
3164 default:
3165 return;
3166 }
3167
3168 ac_nir_build_if(&if_ctx, ctx,
3169 LLVMBuildICmp(ctx->ac.builder, LLVMIntEQ,
3170 invocation_id, ctx->ac.i32_0, ""));
3171
3172 lds_base = get_tcs_out_current_patch_data_offset(ctx);
3173
3174 if (inner_comps) {
3175 tess_inner_index = shader_io_get_unique_index(VARYING_SLOT_TESS_LEVEL_INNER);
3176 lds_inner = LLVMBuildAdd(ctx->ac.builder, lds_base,
3177 LLVMConstInt(ctx->ac.i32, tess_inner_index * 4, false), "");
3178 }
3179
3180 tess_outer_index = shader_io_get_unique_index(VARYING_SLOT_TESS_LEVEL_OUTER);
3181 lds_outer = LLVMBuildAdd(ctx->ac.builder, lds_base,
3182 LLVMConstInt(ctx->ac.i32, tess_outer_index * 4, false), "");
3183
3184 for (i = 0; i < 4; i++) {
3185 inner[i] = LLVMGetUndef(ctx->ac.i32);
3186 outer[i] = LLVMGetUndef(ctx->ac.i32);
3187 }
3188
3189 // LINES reversal
3190 if (ctx->options->key.tcs.primitive_mode == GL_ISOLINES) {
3191 outer[0] = out[1] = ac_lds_load(&ctx->ac, lds_outer);
3192 lds_outer = LLVMBuildAdd(ctx->ac.builder, lds_outer,
3193 ctx->ac.i32_1, "");
3194 outer[1] = out[0] = ac_lds_load(&ctx->ac, lds_outer);
3195 } else {
3196 for (i = 0; i < outer_comps; i++) {
3197 outer[i] = out[i] =
3198 ac_lds_load(&ctx->ac, lds_outer);
3199 lds_outer = LLVMBuildAdd(ctx->ac.builder, lds_outer,
3200 ctx->ac.i32_1, "");
3201 }
3202 for (i = 0; i < inner_comps; i++) {
3203 inner[i] = out[outer_comps+i] =
3204 ac_lds_load(&ctx->ac, lds_inner);
3205 lds_inner = LLVMBuildAdd(ctx->ac.builder, lds_inner,
3206 ctx->ac.i32_1, "");
3207 }
3208 }
3209
3210 /* Convert the outputs to vectors for stores. */
3211 vec0 = ac_build_gather_values(&ctx->ac, out, MIN2(stride, 4));
3212 vec1 = NULL;
3213
3214 if (stride > 4)
3215 vec1 = ac_build_gather_values(&ctx->ac, out + 4, stride - 4);
3216
3217
3218 buffer = ctx->hs_ring_tess_factor;
3219 tf_base = ctx->tess_factor_offset;
3220 byteoffset = LLVMBuildMul(ctx->ac.builder, rel_patch_id,
3221 LLVMConstInt(ctx->ac.i32, 4 * stride, false), "");
3222 unsigned tf_offset = 0;
3223
3224 if (ctx->options->chip_class <= VI) {
3225 ac_nir_build_if(&inner_if_ctx, ctx,
3226 LLVMBuildICmp(ctx->ac.builder, LLVMIntEQ,
3227 rel_patch_id, ctx->ac.i32_0, ""));
3228
3229 /* Store the dynamic HS control word. */
3230 ac_build_buffer_store_dword(&ctx->ac, buffer,
3231 LLVMConstInt(ctx->ac.i32, 0x80000000, false),
3232 1, ctx->ac.i32_0, tf_base,
3233 0, 1, 0, true, false);
3234 tf_offset += 4;
3235
3236 ac_nir_build_endif(&inner_if_ctx);
3237 }
3238
3239 /* Store the tessellation factors. */
3240 ac_build_buffer_store_dword(&ctx->ac, buffer, vec0,
3241 MIN2(stride, 4), byteoffset, tf_base,
3242 tf_offset, 1, 0, true, false);
3243 if (vec1)
3244 ac_build_buffer_store_dword(&ctx->ac, buffer, vec1,
3245 stride - 4, byteoffset, tf_base,
3246 16 + tf_offset, 1, 0, true, false);
3247
3248 //store to offchip for TES to read - only if TES reads them
3249 if (ctx->options->key.tcs.tes_reads_tess_factors) {
3250 LLVMValueRef inner_vec, outer_vec, tf_outer_offset;
3251 LLVMValueRef tf_inner_offset;
3252 unsigned param_outer, param_inner;
3253
3254 param_outer = shader_io_get_unique_index(VARYING_SLOT_TESS_LEVEL_OUTER);
3255 tf_outer_offset = get_tcs_tes_buffer_address(ctx, NULL,
3256 LLVMConstInt(ctx->ac.i32, param_outer, 0));
3257
3258 outer_vec = ac_build_gather_values(&ctx->ac, outer,
3259 util_next_power_of_two(outer_comps));
3260
3261 ac_build_buffer_store_dword(&ctx->ac, ctx->hs_ring_tess_offchip, outer_vec,
3262 outer_comps, tf_outer_offset,
3263 ctx->oc_lds, 0, 1, 0, true, false);
3264 if (inner_comps) {
3265 param_inner = shader_io_get_unique_index(VARYING_SLOT_TESS_LEVEL_INNER);
3266 tf_inner_offset = get_tcs_tes_buffer_address(ctx, NULL,
3267 LLVMConstInt(ctx->ac.i32, param_inner, 0));
3268
3269 inner_vec = inner_comps == 1 ? inner[0] :
3270 ac_build_gather_values(&ctx->ac, inner, inner_comps);
3271 ac_build_buffer_store_dword(&ctx->ac, ctx->hs_ring_tess_offchip, inner_vec,
3272 inner_comps, tf_inner_offset,
3273 ctx->oc_lds, 0, 1, 0, true, false);
3274 }
3275 }
3276 ac_nir_build_endif(&if_ctx);
3277 }
3278
3279 static void
3280 handle_tcs_outputs_post(struct radv_shader_context *ctx)
3281 {
3282 write_tess_factors(ctx);
3283 }
3284
3285 static bool
3286 si_export_mrt_color(struct radv_shader_context *ctx,
3287 LLVMValueRef *color, unsigned index,
3288 struct ac_export_args *args)
3289 {
3290 /* Export */
3291 si_llvm_init_export_args(ctx, color, 0xf,
3292 V_008DFC_SQ_EXP_MRT + index, args);
3293 if (!args->enabled_channels)
3294 return false; /* unnecessary NULL export */
3295
3296 return true;
3297 }
3298
3299 static void
3300 radv_export_mrt_z(struct radv_shader_context *ctx,
3301 LLVMValueRef depth, LLVMValueRef stencil,
3302 LLVMValueRef samplemask)
3303 {
3304 struct ac_export_args args;
3305
3306 ac_export_mrt_z(&ctx->ac, depth, stencil, samplemask, &args);
3307
3308 ac_build_export(&ctx->ac, &args);
3309 }
3310
3311 static void
3312 handle_fs_outputs_post(struct radv_shader_context *ctx)
3313 {
3314 unsigned index = 0;
3315 LLVMValueRef depth = NULL, stencil = NULL, samplemask = NULL;
3316 struct ac_export_args color_args[8];
3317
3318 for (unsigned i = 0; i < AC_LLVM_MAX_OUTPUTS; ++i) {
3319 LLVMValueRef values[4];
3320
3321 if (!(ctx->output_mask & (1ull << i)))
3322 continue;
3323
3324 if (i < FRAG_RESULT_DATA0)
3325 continue;
3326
3327 for (unsigned j = 0; j < 4; j++)
3328 values[j] = ac_to_float(&ctx->ac,
3329 radv_load_output(ctx, i, j));
3330
3331 bool ret = si_export_mrt_color(ctx, values,
3332 i - FRAG_RESULT_DATA0,
3333 &color_args[index]);
3334 if (ret)
3335 index++;
3336 }
3337
3338 /* Process depth, stencil, samplemask. */
3339 if (ctx->shader_info->info.ps.writes_z) {
3340 depth = ac_to_float(&ctx->ac,
3341 radv_load_output(ctx, FRAG_RESULT_DEPTH, 0));
3342 }
3343 if (ctx->shader_info->info.ps.writes_stencil) {
3344 stencil = ac_to_float(&ctx->ac,
3345 radv_load_output(ctx, FRAG_RESULT_STENCIL, 0));
3346 }
3347 if (ctx->shader_info->info.ps.writes_sample_mask) {
3348 samplemask = ac_to_float(&ctx->ac,
3349 radv_load_output(ctx, FRAG_RESULT_SAMPLE_MASK, 0));
3350 }
3351
3352 /* Set the DONE bit on last non-null color export only if Z isn't
3353 * exported.
3354 */
3355 if (index > 0 &&
3356 !ctx->shader_info->info.ps.writes_z &&
3357 !ctx->shader_info->info.ps.writes_stencil &&
3358 !ctx->shader_info->info.ps.writes_sample_mask) {
3359 unsigned last = index - 1;
3360
3361 color_args[last].valid_mask = 1; /* whether the EXEC mask is valid */
3362 color_args[last].done = 1; /* DONE bit */
3363 }
3364
3365 /* Export PS outputs. */
3366 for (unsigned i = 0; i < index; i++)
3367 ac_build_export(&ctx->ac, &color_args[i]);
3368
3369 if (depth || stencil || samplemask)
3370 radv_export_mrt_z(ctx, depth, stencil, samplemask);
3371 else if (!index)
3372 ac_build_export_null(&ctx->ac);
3373 }
3374
3375 static void
3376 emit_gs_epilogue(struct radv_shader_context *ctx)
3377 {
3378 ac_build_sendmsg(&ctx->ac, AC_SENDMSG_GS_OP_NOP | AC_SENDMSG_GS_DONE, ctx->gs_wave_id);
3379 }
3380
3381 static void
3382 handle_shader_outputs_post(struct ac_shader_abi *abi, unsigned max_outputs,
3383 LLVMValueRef *addrs)
3384 {
3385 struct radv_shader_context *ctx = radv_shader_context_from_abi(abi);
3386
3387 switch (ctx->stage) {
3388 case MESA_SHADER_VERTEX:
3389 if (ctx->options->key.vs.as_ls)
3390 handle_ls_outputs_post(ctx);
3391 else if (ctx->options->key.vs.as_es)
3392 handle_es_outputs_post(ctx, &ctx->shader_info->vs.es_info);
3393 else
3394 handle_vs_outputs_post(ctx, ctx->options->key.vs.export_prim_id,
3395 ctx->options->key.vs.export_layer_id,
3396 &ctx->shader_info->vs.outinfo);
3397 break;
3398 case MESA_SHADER_FRAGMENT:
3399 handle_fs_outputs_post(ctx);
3400 break;
3401 case MESA_SHADER_GEOMETRY:
3402 emit_gs_epilogue(ctx);
3403 break;
3404 case MESA_SHADER_TESS_CTRL:
3405 handle_tcs_outputs_post(ctx);
3406 break;
3407 case MESA_SHADER_TESS_EVAL:
3408 if (ctx->options->key.tes.as_es)
3409 handle_es_outputs_post(ctx, &ctx->shader_info->tes.es_info);
3410 else
3411 handle_vs_outputs_post(ctx, ctx->options->key.tes.export_prim_id,
3412 ctx->options->key.tes.export_layer_id,
3413 &ctx->shader_info->tes.outinfo);
3414 break;
3415 default:
3416 break;
3417 }
3418 }
3419
3420 static void ac_llvm_finalize_module(struct radv_shader_context *ctx,
3421 LLVMPassManagerRef passmgr,
3422 const struct radv_nir_compiler_options *options)
3423 {
3424 LLVMRunPassManager(passmgr, ctx->ac.module);
3425 LLVMDisposeBuilder(ctx->ac.builder);
3426
3427 ac_llvm_context_dispose(&ctx->ac);
3428 }
3429
3430 static void
3431 ac_nir_eliminate_const_vs_outputs(struct radv_shader_context *ctx)
3432 {
3433 struct radv_vs_output_info *outinfo;
3434
3435 switch (ctx->stage) {
3436 case MESA_SHADER_FRAGMENT:
3437 case MESA_SHADER_COMPUTE:
3438 case MESA_SHADER_TESS_CTRL:
3439 case MESA_SHADER_GEOMETRY:
3440 return;
3441 case MESA_SHADER_VERTEX:
3442 if (ctx->options->key.vs.as_ls ||
3443 ctx->options->key.vs.as_es)
3444 return;
3445 outinfo = &ctx->shader_info->vs.outinfo;
3446 break;
3447 case MESA_SHADER_TESS_EVAL:
3448 if (ctx->options->key.vs.as_es)
3449 return;
3450 outinfo = &ctx->shader_info->tes.outinfo;
3451 break;
3452 default:
3453 unreachable("Unhandled shader type");
3454 }
3455
3456 ac_optimize_vs_outputs(&ctx->ac,
3457 ctx->main_function,
3458 outinfo->vs_output_param_offset,
3459 VARYING_SLOT_MAX,
3460 &outinfo->param_exports);
3461 }
3462
3463 static void
3464 ac_setup_rings(struct radv_shader_context *ctx)
3465 {
3466 if (ctx->options->chip_class <= VI &&
3467 (ctx->stage == MESA_SHADER_GEOMETRY ||
3468 ctx->options->key.vs.as_es || ctx->options->key.tes.as_es)) {
3469 unsigned ring = ctx->stage == MESA_SHADER_GEOMETRY ? RING_ESGS_GS
3470 : RING_ESGS_VS;
3471 LLVMValueRef offset = LLVMConstInt(ctx->ac.i32, ring, false);
3472
3473 ctx->esgs_ring = ac_build_load_to_sgpr(&ctx->ac,
3474 ctx->ring_offsets,
3475 offset);
3476 }
3477
3478 if (ctx->is_gs_copy_shader) {
3479 ctx->gsvs_ring[0] =
3480 ac_build_load_to_sgpr(&ctx->ac, ctx->ring_offsets,
3481 LLVMConstInt(ctx->ac.i32,
3482 RING_GSVS_VS, false));
3483 }
3484
3485 if (ctx->stage == MESA_SHADER_GEOMETRY) {
3486 /* The conceptual layout of the GSVS ring is
3487 * v0c0 .. vLv0 v0c1 .. vLc1 ..
3488 * but the real memory layout is swizzled across
3489 * threads:
3490 * t0v0c0 .. t15v0c0 t0v1c0 .. t15v1c0 ... t15vLcL
3491 * t16v0c0 ..
3492 * Override the buffer descriptor accordingly.
3493 */
3494 LLVMTypeRef v2i64 = LLVMVectorType(ctx->ac.i64, 2);
3495 uint64_t stream_offset = 0;
3496 unsigned num_records = 64;
3497 LLVMValueRef base_ring;
3498
3499 base_ring =
3500 ac_build_load_to_sgpr(&ctx->ac, ctx->ring_offsets,
3501 LLVMConstInt(ctx->ac.i32,
3502 RING_GSVS_GS, false));
3503
3504 for (unsigned stream = 0; stream < 4; stream++) {
3505 unsigned num_components, stride;
3506 LLVMValueRef ring, tmp;
3507
3508 num_components =
3509 ctx->shader_info->info.gs.num_stream_output_components[stream];
3510
3511 if (!num_components)
3512 continue;
3513
3514 stride = 4 * num_components * ctx->gs_max_out_vertices;
3515
3516 /* Limit on the stride field for <= CIK. */
3517 assert(stride < (1 << 14));
3518
3519 ring = LLVMBuildBitCast(ctx->ac.builder,
3520 base_ring, v2i64, "");
3521 tmp = LLVMBuildExtractElement(ctx->ac.builder,
3522 ring, ctx->ac.i32_0, "");
3523 tmp = LLVMBuildAdd(ctx->ac.builder, tmp,
3524 LLVMConstInt(ctx->ac.i64,
3525 stream_offset, 0), "");
3526 ring = LLVMBuildInsertElement(ctx->ac.builder,
3527 ring, tmp, ctx->ac.i32_0, "");
3528
3529 stream_offset += stride * 64;
3530
3531 ring = LLVMBuildBitCast(ctx->ac.builder, ring,
3532 ctx->ac.v4i32, "");
3533
3534 tmp = LLVMBuildExtractElement(ctx->ac.builder, ring,
3535 ctx->ac.i32_1, "");
3536 tmp = LLVMBuildOr(ctx->ac.builder, tmp,
3537 LLVMConstInt(ctx->ac.i32,
3538 S_008F04_STRIDE(stride), false), "");
3539 ring = LLVMBuildInsertElement(ctx->ac.builder, ring, tmp,
3540 ctx->ac.i32_1, "");
3541
3542 ring = LLVMBuildInsertElement(ctx->ac.builder, ring,
3543 LLVMConstInt(ctx->ac.i32,
3544 num_records, false),
3545 LLVMConstInt(ctx->ac.i32, 2, false), "");
3546
3547 ctx->gsvs_ring[stream] = ring;
3548 }
3549 }
3550
3551 if (ctx->stage == MESA_SHADER_TESS_CTRL ||
3552 ctx->stage == MESA_SHADER_TESS_EVAL) {
3553 ctx->hs_ring_tess_offchip = ac_build_load_to_sgpr(&ctx->ac, ctx->ring_offsets, LLVMConstInt(ctx->ac.i32, RING_HS_TESS_OFFCHIP, false));
3554 ctx->hs_ring_tess_factor = ac_build_load_to_sgpr(&ctx->ac, ctx->ring_offsets, LLVMConstInt(ctx->ac.i32, RING_HS_TESS_FACTOR, false));
3555 }
3556 }
3557
3558 unsigned
3559 radv_nir_get_max_workgroup_size(enum chip_class chip_class,
3560 const struct nir_shader *nir)
3561 {
3562 switch (nir->info.stage) {
3563 case MESA_SHADER_TESS_CTRL:
3564 return chip_class >= CIK ? 128 : 64;
3565 case MESA_SHADER_GEOMETRY:
3566 return chip_class >= GFX9 ? 128 : 64;
3567 case MESA_SHADER_COMPUTE:
3568 break;
3569 default:
3570 return 0;
3571 }
3572
3573 unsigned max_workgroup_size = nir->info.cs.local_size[0] *
3574 nir->info.cs.local_size[1] *
3575 nir->info.cs.local_size[2];
3576 return max_workgroup_size;
3577 }
3578
3579 /* Fixup the HW not emitting the TCS regs if there are no HS threads. */
3580 static void ac_nir_fixup_ls_hs_input_vgprs(struct radv_shader_context *ctx)
3581 {
3582 LLVMValueRef count = ac_unpack_param(&ctx->ac, ctx->merged_wave_info, 8, 8);
3583 LLVMValueRef hs_empty = LLVMBuildICmp(ctx->ac.builder, LLVMIntEQ, count,
3584 ctx->ac.i32_0, "");
3585 ctx->abi.instance_id = LLVMBuildSelect(ctx->ac.builder, hs_empty, ctx->rel_auto_id, ctx->abi.instance_id, "");
3586 ctx->rel_auto_id = LLVMBuildSelect(ctx->ac.builder, hs_empty, ctx->abi.tcs_rel_ids, ctx->rel_auto_id, "");
3587 ctx->abi.vertex_id = LLVMBuildSelect(ctx->ac.builder, hs_empty, ctx->abi.tcs_patch_id, ctx->abi.vertex_id, "");
3588 }
3589
3590 static void prepare_gs_input_vgprs(struct radv_shader_context *ctx)
3591 {
3592 for(int i = 5; i >= 0; --i) {
3593 ctx->gs_vtx_offset[i] = ac_unpack_param(&ctx->ac, ctx->gs_vtx_offset[i & ~1],
3594 (i & 1) * 16, 16);
3595 }
3596
3597 ctx->gs_wave_id = ac_unpack_param(&ctx->ac, ctx->merged_wave_info, 16, 8);
3598 }
3599
3600
3601 static
3602 LLVMModuleRef ac_translate_nir_to_llvm(struct ac_llvm_compiler *ac_llvm,
3603 struct nir_shader *const *shaders,
3604 int shader_count,
3605 struct radv_shader_variant_info *shader_info,
3606 const struct radv_nir_compiler_options *options)
3607 {
3608 struct radv_shader_context ctx = {0};
3609 unsigned i;
3610 ctx.options = options;
3611 ctx.shader_info = shader_info;
3612
3613 ac_llvm_context_init(&ctx.ac, options->chip_class, options->family);
3614 ctx.context = ctx.ac.context;
3615 ctx.ac.module = ac_create_module(ac_llvm->tm, ctx.context);
3616
3617 enum ac_float_mode float_mode =
3618 options->unsafe_math ? AC_FLOAT_MODE_UNSAFE_FP_MATH :
3619 AC_FLOAT_MODE_DEFAULT;
3620
3621 ctx.ac.builder = ac_create_builder(ctx.context, float_mode);
3622
3623 memset(shader_info, 0, sizeof(*shader_info));
3624
3625 radv_nir_shader_info_init(&shader_info->info);
3626
3627 for(int i = 0; i < shader_count; ++i)
3628 radv_nir_shader_info_pass(shaders[i], options, &shader_info->info);
3629
3630 for (i = 0; i < RADV_UD_MAX_SETS; i++)
3631 shader_info->user_sgprs_locs.descriptor_sets[i].sgpr_idx = -1;
3632 for (i = 0; i < AC_UD_MAX_UD; i++)
3633 shader_info->user_sgprs_locs.shader_data[i].sgpr_idx = -1;
3634
3635 ctx.max_workgroup_size = 0;
3636 for (int i = 0; i < shader_count; ++i) {
3637 ctx.max_workgroup_size = MAX2(ctx.max_workgroup_size,
3638 radv_nir_get_max_workgroup_size(ctx.options->chip_class,
3639 shaders[i]));
3640 }
3641
3642 create_function(&ctx, shaders[shader_count - 1]->info.stage, shader_count >= 2,
3643 shader_count >= 2 ? shaders[shader_count - 2]->info.stage : MESA_SHADER_VERTEX);
3644
3645 ctx.abi.inputs = &ctx.inputs[0];
3646 ctx.abi.emit_outputs = handle_shader_outputs_post;
3647 ctx.abi.emit_vertex = visit_emit_vertex;
3648 ctx.abi.load_ubo = radv_load_ubo;
3649 ctx.abi.load_ssbo = radv_load_ssbo;
3650 ctx.abi.load_sampler_desc = radv_get_sampler_desc;
3651 ctx.abi.load_resource = radv_load_resource;
3652 ctx.abi.clamp_shadow_reference = false;
3653 ctx.abi.gfx9_stride_size_workaround = ctx.ac.chip_class == GFX9 && HAVE_LLVM < 0x800;
3654
3655 if (shader_count >= 2)
3656 ac_init_exec_full_mask(&ctx.ac);
3657
3658 if (ctx.ac.chip_class == GFX9 &&
3659 shaders[shader_count - 1]->info.stage == MESA_SHADER_TESS_CTRL)
3660 ac_nir_fixup_ls_hs_input_vgprs(&ctx);
3661
3662 for(int i = 0; i < shader_count; ++i) {
3663 ctx.stage = shaders[i]->info.stage;
3664 ctx.output_mask = 0;
3665
3666 if (shaders[i]->info.stage == MESA_SHADER_GEOMETRY) {
3667 for (int i = 0; i < 4; i++) {
3668 ctx.gs_next_vertex[i] =
3669 ac_build_alloca(&ctx.ac, ctx.ac.i32, "");
3670 }
3671 ctx.gs_max_out_vertices = shaders[i]->info.gs.vertices_out;
3672 ctx.abi.load_inputs = load_gs_input;
3673 ctx.abi.emit_primitive = visit_end_primitive;
3674 } else if (shaders[i]->info.stage == MESA_SHADER_TESS_CTRL) {
3675 ctx.tcs_outputs_read = shaders[i]->info.outputs_read;
3676 ctx.tcs_patch_outputs_read = shaders[i]->info.patch_outputs_read;
3677 ctx.abi.load_tess_varyings = load_tcs_varyings;
3678 ctx.abi.load_patch_vertices_in = load_patch_vertices_in;
3679 ctx.abi.store_tcs_outputs = store_tcs_output;
3680 ctx.tcs_vertices_per_patch = shaders[i]->info.tess.tcs_vertices_out;
3681 if (shader_count == 1)
3682 ctx.tcs_num_inputs = ctx.options->key.tcs.num_inputs;
3683 else
3684 ctx.tcs_num_inputs = util_last_bit64(shader_info->info.vs.ls_outputs_written);
3685 ctx.tcs_num_patches = get_tcs_num_patches(&ctx);
3686 } else if (shaders[i]->info.stage == MESA_SHADER_TESS_EVAL) {
3687 ctx.tes_primitive_mode = shaders[i]->info.tess.primitive_mode;
3688 ctx.abi.load_tess_varyings = load_tes_input;
3689 ctx.abi.load_tess_coord = load_tess_coord;
3690 ctx.abi.load_patch_vertices_in = load_patch_vertices_in;
3691 ctx.tcs_vertices_per_patch = shaders[i]->info.tess.tcs_vertices_out;
3692 ctx.tcs_num_patches = ctx.options->key.tes.num_patches;
3693 } else if (shaders[i]->info.stage == MESA_SHADER_VERTEX) {
3694 if (shader_info->info.vs.needs_instance_id) {
3695 if (ctx.options->key.vs.as_ls) {
3696 ctx.shader_info->vs.vgpr_comp_cnt =
3697 MAX2(2, ctx.shader_info->vs.vgpr_comp_cnt);
3698 } else {
3699 ctx.shader_info->vs.vgpr_comp_cnt =
3700 MAX2(1, ctx.shader_info->vs.vgpr_comp_cnt);
3701 }
3702 }
3703 ctx.abi.load_base_vertex = radv_load_base_vertex;
3704 } else if (shaders[i]->info.stage == MESA_SHADER_FRAGMENT) {
3705 shader_info->fs.can_discard = shaders[i]->info.fs.uses_discard;
3706 ctx.abi.lookup_interp_param = lookup_interp_param;
3707 ctx.abi.load_sample_position = load_sample_position;
3708 ctx.abi.load_sample_mask_in = load_sample_mask_in;
3709 ctx.abi.emit_kill = radv_emit_kill;
3710 }
3711
3712 if (i)
3713 ac_emit_barrier(&ctx.ac, ctx.stage);
3714
3715 nir_foreach_variable(variable, &shaders[i]->outputs)
3716 scan_shader_output_decl(&ctx, variable, shaders[i], shaders[i]->info.stage);
3717
3718 if (shaders[i]->info.stage == MESA_SHADER_GEOMETRY) {
3719 unsigned addclip = shaders[i]->info.clip_distance_array_size +
3720 shaders[i]->info.cull_distance_array_size > 4;
3721 ctx.gsvs_vertex_size = (util_bitcount64(ctx.output_mask) + addclip) * 16;
3722 ctx.max_gsvs_emit_size = ctx.gsvs_vertex_size *
3723 shaders[i]->info.gs.vertices_out;
3724 }
3725
3726 ac_setup_rings(&ctx);
3727
3728 LLVMBasicBlockRef merge_block;
3729 if (shader_count >= 2) {
3730 LLVMValueRef fn = LLVMGetBasicBlockParent(LLVMGetInsertBlock(ctx.ac.builder));
3731 LLVMBasicBlockRef then_block = LLVMAppendBasicBlockInContext(ctx.ac.context, fn, "");
3732 merge_block = LLVMAppendBasicBlockInContext(ctx.ac.context, fn, "");
3733
3734 LLVMValueRef count = ac_unpack_param(&ctx.ac, ctx.merged_wave_info, 8 * i, 8);
3735 LLVMValueRef thread_id = ac_get_thread_id(&ctx.ac);
3736 LLVMValueRef cond = LLVMBuildICmp(ctx.ac.builder, LLVMIntULT,
3737 thread_id, count, "");
3738 LLVMBuildCondBr(ctx.ac.builder, cond, then_block, merge_block);
3739
3740 LLVMPositionBuilderAtEnd(ctx.ac.builder, then_block);
3741 }
3742
3743 if (shaders[i]->info.stage == MESA_SHADER_FRAGMENT)
3744 handle_fs_inputs(&ctx, shaders[i]);
3745 else if(shaders[i]->info.stage == MESA_SHADER_VERTEX)
3746 handle_vs_inputs(&ctx, shaders[i]);
3747 else if(shader_count >= 2 && shaders[i]->info.stage == MESA_SHADER_GEOMETRY)
3748 prepare_gs_input_vgprs(&ctx);
3749
3750 ac_nir_translate(&ctx.ac, &ctx.abi, shaders[i]);
3751
3752 if (shader_count >= 2) {
3753 LLVMBuildBr(ctx.ac.builder, merge_block);
3754 LLVMPositionBuilderAtEnd(ctx.ac.builder, merge_block);
3755 }
3756
3757 if (shaders[i]->info.stage == MESA_SHADER_GEOMETRY) {
3758 shader_info->gs.gsvs_vertex_size = ctx.gsvs_vertex_size;
3759 shader_info->gs.max_gsvs_emit_size = ctx.max_gsvs_emit_size;
3760 } else if (shaders[i]->info.stage == MESA_SHADER_TESS_CTRL) {
3761 shader_info->tcs.num_patches = ctx.tcs_num_patches;
3762 shader_info->tcs.lds_size = calculate_tess_lds_size(&ctx);
3763 }
3764 }
3765
3766 LLVMBuildRetVoid(ctx.ac.builder);
3767
3768 if (options->dump_preoptir)
3769 ac_dump_module(ctx.ac.module);
3770
3771 ac_llvm_finalize_module(&ctx, ac_llvm->passmgr, options);
3772
3773 if (shader_count == 1)
3774 ac_nir_eliminate_const_vs_outputs(&ctx);
3775
3776 if (options->dump_shader) {
3777 ctx.shader_info->private_mem_vgprs =
3778 ac_count_scratch_private_memory(ctx.main_function);
3779 }
3780
3781 return ctx.ac.module;
3782 }
3783
3784 static void ac_diagnostic_handler(LLVMDiagnosticInfoRef di, void *context)
3785 {
3786 unsigned *retval = (unsigned *)context;
3787 LLVMDiagnosticSeverity severity = LLVMGetDiagInfoSeverity(di);
3788 char *description = LLVMGetDiagInfoDescription(di);
3789
3790 if (severity == LLVMDSError) {
3791 *retval = 1;
3792 fprintf(stderr, "LLVM triggered Diagnostic Handler: %s\n",
3793 description);
3794 }
3795
3796 LLVMDisposeMessage(description);
3797 }
3798
3799 static unsigned ac_llvm_compile(LLVMModuleRef M,
3800 struct ac_shader_binary *binary,
3801 struct ac_llvm_compiler *ac_llvm)
3802 {
3803 unsigned retval = 0;
3804 LLVMContextRef llvm_ctx;
3805
3806 /* Setup Diagnostic Handler*/
3807 llvm_ctx = LLVMGetModuleContext(M);
3808
3809 LLVMContextSetDiagnosticHandler(llvm_ctx, ac_diagnostic_handler,
3810 &retval);
3811
3812 /* Compile IR*/
3813 if (!radv_compile_to_binary(ac_llvm, M, binary))
3814 retval = 1;
3815 return retval;
3816 }
3817
3818 static void ac_compile_llvm_module(struct ac_llvm_compiler *ac_llvm,
3819 LLVMModuleRef llvm_module,
3820 struct ac_shader_binary *binary,
3821 struct ac_shader_config *config,
3822 struct radv_shader_variant_info *shader_info,
3823 gl_shader_stage stage,
3824 const struct radv_nir_compiler_options *options)
3825 {
3826 if (options->dump_shader)
3827 ac_dump_module(llvm_module);
3828
3829 memset(binary, 0, sizeof(*binary));
3830
3831 if (options->record_llvm_ir) {
3832 char *llvm_ir = LLVMPrintModuleToString(llvm_module);
3833 binary->llvm_ir_string = strdup(llvm_ir);
3834 LLVMDisposeMessage(llvm_ir);
3835 }
3836
3837 int v = ac_llvm_compile(llvm_module, binary, ac_llvm);
3838 if (v) {
3839 fprintf(stderr, "compile failed\n");
3840 }
3841
3842 if (options->dump_shader)
3843 fprintf(stderr, "disasm:\n%s\n", binary->disasm_string);
3844
3845 ac_shader_binary_read_config(binary, config, 0, options->supports_spill);
3846
3847 LLVMContextRef ctx = LLVMGetModuleContext(llvm_module);
3848 LLVMDisposeModule(llvm_module);
3849 LLVMContextDispose(ctx);
3850
3851 if (stage == MESA_SHADER_FRAGMENT) {
3852 shader_info->num_input_vgprs = 0;
3853 if (G_0286CC_PERSP_SAMPLE_ENA(config->spi_ps_input_addr))
3854 shader_info->num_input_vgprs += 2;
3855 if (G_0286CC_PERSP_CENTER_ENA(config->spi_ps_input_addr))
3856 shader_info->num_input_vgprs += 2;
3857 if (G_0286CC_PERSP_CENTROID_ENA(config->spi_ps_input_addr))
3858 shader_info->num_input_vgprs += 2;
3859 if (G_0286CC_PERSP_PULL_MODEL_ENA(config->spi_ps_input_addr))
3860 shader_info->num_input_vgprs += 3;
3861 if (G_0286CC_LINEAR_SAMPLE_ENA(config->spi_ps_input_addr))
3862 shader_info->num_input_vgprs += 2;
3863 if (G_0286CC_LINEAR_CENTER_ENA(config->spi_ps_input_addr))
3864 shader_info->num_input_vgprs += 2;
3865 if (G_0286CC_LINEAR_CENTROID_ENA(config->spi_ps_input_addr))
3866 shader_info->num_input_vgprs += 2;
3867 if (G_0286CC_LINE_STIPPLE_TEX_ENA(config->spi_ps_input_addr))
3868 shader_info->num_input_vgprs += 1;
3869 if (G_0286CC_POS_X_FLOAT_ENA(config->spi_ps_input_addr))
3870 shader_info->num_input_vgprs += 1;
3871 if (G_0286CC_POS_Y_FLOAT_ENA(config->spi_ps_input_addr))
3872 shader_info->num_input_vgprs += 1;
3873 if (G_0286CC_POS_Z_FLOAT_ENA(config->spi_ps_input_addr))
3874 shader_info->num_input_vgprs += 1;
3875 if (G_0286CC_POS_W_FLOAT_ENA(config->spi_ps_input_addr))
3876 shader_info->num_input_vgprs += 1;
3877 if (G_0286CC_FRONT_FACE_ENA(config->spi_ps_input_addr))
3878 shader_info->num_input_vgprs += 1;
3879 if (G_0286CC_ANCILLARY_ENA(config->spi_ps_input_addr))
3880 shader_info->num_input_vgprs += 1;
3881 if (G_0286CC_SAMPLE_COVERAGE_ENA(config->spi_ps_input_addr))
3882 shader_info->num_input_vgprs += 1;
3883 if (G_0286CC_POS_FIXED_PT_ENA(config->spi_ps_input_addr))
3884 shader_info->num_input_vgprs += 1;
3885 }
3886 config->num_vgprs = MAX2(config->num_vgprs, shader_info->num_input_vgprs);
3887
3888 /* +3 for scratch wave offset and VCC */
3889 config->num_sgprs = MAX2(config->num_sgprs,
3890 shader_info->num_input_sgprs + 3);
3891
3892 /* Enable 64-bit and 16-bit denormals, because there is no performance
3893 * cost.
3894 *
3895 * If denormals are enabled, all floating-point output modifiers are
3896 * ignored.
3897 *
3898 * Don't enable denormals for 32-bit floats, because:
3899 * - Floating-point output modifiers would be ignored by the hw.
3900 * - Some opcodes don't support denormals, such as v_mad_f32. We would
3901 * have to stop using those.
3902 * - SI & CI would be very slow.
3903 */
3904 config->float_mode |= V_00B028_FP_64_DENORMS;
3905 }
3906
3907 static void
3908 ac_fill_shader_info(struct radv_shader_variant_info *shader_info, struct nir_shader *nir, const struct radv_nir_compiler_options *options)
3909 {
3910 switch (nir->info.stage) {
3911 case MESA_SHADER_COMPUTE:
3912 for (int i = 0; i < 3; ++i)
3913 shader_info->cs.block_size[i] = nir->info.cs.local_size[i];
3914 break;
3915 case MESA_SHADER_FRAGMENT:
3916 shader_info->fs.early_fragment_test = nir->info.fs.early_fragment_tests;
3917 break;
3918 case MESA_SHADER_GEOMETRY:
3919 shader_info->gs.vertices_in = nir->info.gs.vertices_in;
3920 shader_info->gs.vertices_out = nir->info.gs.vertices_out;
3921 shader_info->gs.output_prim = nir->info.gs.output_primitive;
3922 shader_info->gs.invocations = nir->info.gs.invocations;
3923 break;
3924 case MESA_SHADER_TESS_EVAL:
3925 shader_info->tes.primitive_mode = nir->info.tess.primitive_mode;
3926 shader_info->tes.spacing = nir->info.tess.spacing;
3927 shader_info->tes.ccw = nir->info.tess.ccw;
3928 shader_info->tes.point_mode = nir->info.tess.point_mode;
3929 shader_info->tes.as_es = options->key.tes.as_es;
3930 break;
3931 case MESA_SHADER_TESS_CTRL:
3932 shader_info->tcs.tcs_vertices_out = nir->info.tess.tcs_vertices_out;
3933 break;
3934 case MESA_SHADER_VERTEX:
3935 shader_info->vs.as_es = options->key.vs.as_es;
3936 shader_info->vs.as_ls = options->key.vs.as_ls;
3937 /* in LS mode we need at least 1, invocation id needs 2, handled elsewhere */
3938 if (options->key.vs.as_ls)
3939 shader_info->vs.vgpr_comp_cnt = MAX2(1, shader_info->vs.vgpr_comp_cnt);
3940 break;
3941 default:
3942 break;
3943 }
3944 }
3945
3946 void
3947 radv_compile_nir_shader(struct ac_llvm_compiler *ac_llvm,
3948 struct ac_shader_binary *binary,
3949 struct ac_shader_config *config,
3950 struct radv_shader_variant_info *shader_info,
3951 struct nir_shader *const *nir,
3952 int nir_count,
3953 const struct radv_nir_compiler_options *options)
3954 {
3955
3956 LLVMModuleRef llvm_module;
3957
3958 llvm_module = ac_translate_nir_to_llvm(ac_llvm, nir, nir_count, shader_info,
3959 options);
3960
3961 ac_compile_llvm_module(ac_llvm, llvm_module, binary, config, shader_info,
3962 nir[0]->info.stage, options);
3963
3964 for (int i = 0; i < nir_count; ++i)
3965 ac_fill_shader_info(shader_info, nir[i], options);
3966
3967 /* Determine the ES type (VS or TES) for the GS on GFX9. */
3968 if (options->chip_class == GFX9) {
3969 if (nir_count == 2 &&
3970 nir[1]->info.stage == MESA_SHADER_GEOMETRY) {
3971 shader_info->gs.es_type = nir[0]->info.stage;
3972 }
3973 }
3974 }
3975
3976 static void
3977 ac_gs_copy_shader_emit(struct radv_shader_context *ctx)
3978 {
3979 LLVMValueRef vtx_offset =
3980 LLVMBuildMul(ctx->ac.builder, ctx->abi.vertex_id,
3981 LLVMConstInt(ctx->ac.i32, 4, false), "");
3982 LLVMValueRef stream_id;
3983
3984 /* Fetch the vertex stream ID. */
3985 if (ctx->shader_info->info.so.num_outputs) {
3986 stream_id =
3987 ac_unpack_param(&ctx->ac, ctx->streamout_config, 24, 2);
3988 } else {
3989 stream_id = ctx->ac.i32_0;
3990 }
3991
3992 LLVMBasicBlockRef end_bb;
3993 LLVMValueRef switch_inst;
3994
3995 end_bb = LLVMAppendBasicBlockInContext(ctx->ac.context,
3996 ctx->main_function, "end");
3997 switch_inst = LLVMBuildSwitch(ctx->ac.builder, stream_id, end_bb, 4);
3998
3999 for (unsigned stream = 0; stream < 4; stream++) {
4000 unsigned num_components =
4001 ctx->shader_info->info.gs.num_stream_output_components[stream];
4002 LLVMBasicBlockRef bb;
4003 unsigned offset;
4004
4005 if (!num_components)
4006 continue;
4007
4008 if (stream > 0 && !ctx->shader_info->info.so.num_outputs)
4009 continue;
4010
4011 bb = LLVMInsertBasicBlockInContext(ctx->ac.context, end_bb, "out");
4012 LLVMAddCase(switch_inst, LLVMConstInt(ctx->ac.i32, stream, 0), bb);
4013 LLVMPositionBuilderAtEnd(ctx->ac.builder, bb);
4014
4015 offset = 0;
4016 for (unsigned i = 0; i < AC_LLVM_MAX_OUTPUTS; ++i) {
4017 unsigned output_usage_mask =
4018 ctx->shader_info->info.gs.output_usage_mask[i];
4019 unsigned output_stream =
4020 ctx->shader_info->info.gs.output_streams[i];
4021 int length = util_last_bit(output_usage_mask);
4022
4023 if (!(ctx->output_mask & (1ull << i)) ||
4024 output_stream != stream)
4025 continue;
4026
4027 for (unsigned j = 0; j < length; j++) {
4028 LLVMValueRef value, soffset;
4029
4030 if (!(output_usage_mask & (1 << j)))
4031 continue;
4032
4033 soffset = LLVMConstInt(ctx->ac.i32,
4034 offset *
4035 ctx->gs_max_out_vertices * 16 * 4, false);
4036
4037 offset++;
4038
4039 value = ac_build_buffer_load(&ctx->ac,
4040 ctx->gsvs_ring[0],
4041 1, ctx->ac.i32_0,
4042 vtx_offset, soffset,
4043 0, 1, 1, true, false);
4044
4045 LLVMTypeRef type = LLVMGetAllocatedType(ctx->abi.outputs[ac_llvm_reg_index_soa(i, j)]);
4046 if (ac_get_type_size(type) == 2) {
4047 value = LLVMBuildBitCast(ctx->ac.builder, value, ctx->ac.i32, "");
4048 value = LLVMBuildTrunc(ctx->ac.builder, value, ctx->ac.i16, "");
4049 }
4050
4051 LLVMBuildStore(ctx->ac.builder,
4052 ac_to_float(&ctx->ac, value), ctx->abi.outputs[ac_llvm_reg_index_soa(i, j)]);
4053 }
4054 }
4055
4056 if (ctx->shader_info->info.so.num_outputs)
4057 radv_emit_streamout(ctx, stream);
4058
4059 if (stream == 0) {
4060 handle_vs_outputs_post(ctx, false, false,
4061 &ctx->shader_info->vs.outinfo);
4062 }
4063
4064 LLVMBuildBr(ctx->ac.builder, end_bb);
4065 }
4066
4067 LLVMPositionBuilderAtEnd(ctx->ac.builder, end_bb);
4068 }
4069
4070 void
4071 radv_compile_gs_copy_shader(struct ac_llvm_compiler *ac_llvm,
4072 struct nir_shader *geom_shader,
4073 struct ac_shader_binary *binary,
4074 struct ac_shader_config *config,
4075 struct radv_shader_variant_info *shader_info,
4076 const struct radv_nir_compiler_options *options)
4077 {
4078 struct radv_shader_context ctx = {0};
4079 ctx.options = options;
4080 ctx.shader_info = shader_info;
4081
4082 ac_llvm_context_init(&ctx.ac, options->chip_class, options->family);
4083 ctx.context = ctx.ac.context;
4084 ctx.ac.module = ac_create_module(ac_llvm->tm, ctx.context);
4085
4086 ctx.is_gs_copy_shader = true;
4087
4088 enum ac_float_mode float_mode =
4089 options->unsafe_math ? AC_FLOAT_MODE_UNSAFE_FP_MATH :
4090 AC_FLOAT_MODE_DEFAULT;
4091
4092 ctx.ac.builder = ac_create_builder(ctx.context, float_mode);
4093 ctx.stage = MESA_SHADER_VERTEX;
4094
4095 radv_nir_shader_info_pass(geom_shader, options, &shader_info->info);
4096
4097 create_function(&ctx, MESA_SHADER_VERTEX, false, MESA_SHADER_VERTEX);
4098
4099 ctx.gs_max_out_vertices = geom_shader->info.gs.vertices_out;
4100 ac_setup_rings(&ctx);
4101
4102 nir_foreach_variable(variable, &geom_shader->outputs) {
4103 scan_shader_output_decl(&ctx, variable, geom_shader, MESA_SHADER_VERTEX);
4104 ac_handle_shader_output_decl(&ctx.ac, &ctx.abi, geom_shader,
4105 variable, MESA_SHADER_VERTEX);
4106 }
4107
4108 ac_gs_copy_shader_emit(&ctx);
4109
4110 LLVMBuildRetVoid(ctx.ac.builder);
4111
4112 ac_llvm_finalize_module(&ctx, ac_llvm->passmgr, options);
4113
4114 ac_compile_llvm_module(ac_llvm, ctx.ac.module, binary, config, shader_info,
4115 MESA_SHADER_VERTEX, options);
4116 }