projects / mesa.git / blob
? search:
summary | shortlog | log | commit | commitdiff | tree
history | raw | HEAD
radv: apply the indexing workaround for atomic buffer operations on GFX9
[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 = LLVMConstInt(ctx->ac.i32, base_offset, false);
1309
1310 if (layout->binding[binding].type != VK_DESCRIPTOR_TYPE_INLINE_UNIFORM_BLOCK_EXT) {
1311 offset = ac_build_imad(&ctx->ac, index, stride, offset);
1312 }
1313
1314 desc_ptr = ac_build_gep0(&ctx->ac, desc_ptr, offset);
1315 desc_ptr = ac_cast_ptr(&ctx->ac, desc_ptr, ctx->ac.v4i32);
1316 LLVMSetMetadata(desc_ptr, ctx->ac.uniform_md_kind, ctx->ac.empty_md);
1317
1318 if (layout->binding[binding].type == VK_DESCRIPTOR_TYPE_INLINE_UNIFORM_BLOCK_EXT) {
1319 uint32_t desc_type = S_008F0C_DST_SEL_X(V_008F0C_SQ_SEL_X) |
1320 S_008F0C_DST_SEL_Y(V_008F0C_SQ_SEL_Y) |
1321 S_008F0C_DST_SEL_Z(V_008F0C_SQ_SEL_Z) |
1322 S_008F0C_DST_SEL_W(V_008F0C_SQ_SEL_W) |
1323 S_008F0C_NUM_FORMAT(V_008F0C_BUF_NUM_FORMAT_FLOAT) |
1324 S_008F0C_DATA_FORMAT(V_008F0C_BUF_DATA_FORMAT_32);
1325
1326 LLVMValueRef desc_components[4] = {
1327 LLVMBuildPtrToInt(ctx->ac.builder, desc_ptr, ctx->ac.intptr, ""),
1328 LLVMConstInt(ctx->ac.i32, S_008F04_BASE_ADDRESS_HI(ctx->options->address32_hi), false),
1329 /* High limit to support variable sizes. */
1330 LLVMConstInt(ctx->ac.i32, 0xffffffff, false),
1331 LLVMConstInt(ctx->ac.i32, desc_type, false),
1332 };
1333
1334 return ac_build_gather_values(&ctx->ac, desc_components, 4);
1335 }
1336
1337 return desc_ptr;
1338 }
1339
1340
1341 /* The offchip buffer layout for TCS->TES is
1342 *
1343 * - attribute 0 of patch 0 vertex 0
1344 * - attribute 0 of patch 0 vertex 1
1345 * - attribute 0 of patch 0 vertex 2
1346 * ...
1347 * - attribute 0 of patch 1 vertex 0
1348 * - attribute 0 of patch 1 vertex 1
1349 * ...
1350 * - attribute 1 of patch 0 vertex 0
1351 * - attribute 1 of patch 0 vertex 1
1352 * ...
1353 * - per patch attribute 0 of patch 0
1354 * - per patch attribute 0 of patch 1
1355 * ...
1356 *
1357 * Note that every attribute has 4 components.
1358 */
1359 static LLVMValueRef get_non_vertex_index_offset(struct radv_shader_context *ctx)
1360 {
1361 uint32_t num_patches = ctx->tcs_num_patches;
1362 uint32_t num_tcs_outputs;
1363 if (ctx->stage == MESA_SHADER_TESS_CTRL)
1364 num_tcs_outputs = util_last_bit64(ctx->shader_info->info.tcs.outputs_written);
1365 else
1366 num_tcs_outputs = ctx->options->key.tes.tcs_num_outputs;
1367
1368 uint32_t output_vertex_size = num_tcs_outputs * 16;
1369 uint32_t pervertex_output_patch_size = ctx->tcs_vertices_per_patch * output_vertex_size;
1370
1371 return LLVMConstInt(ctx->ac.i32, pervertex_output_patch_size * num_patches, false);
1372 }
1373
1374 static LLVMValueRef calc_param_stride(struct radv_shader_context *ctx,
1375 LLVMValueRef vertex_index)
1376 {
1377 LLVMValueRef param_stride;
1378 if (vertex_index)
1379 param_stride = LLVMConstInt(ctx->ac.i32, ctx->tcs_vertices_per_patch * ctx->tcs_num_patches, false);
1380 else
1381 param_stride = LLVMConstInt(ctx->ac.i32, ctx->tcs_num_patches, false);
1382 return param_stride;
1383 }
1384
1385 static LLVMValueRef get_tcs_tes_buffer_address(struct radv_shader_context *ctx,
1386 LLVMValueRef vertex_index,
1387 LLVMValueRef param_index)
1388 {
1389 LLVMValueRef base_addr;
1390 LLVMValueRef param_stride, constant16;
1391 LLVMValueRef rel_patch_id = get_rel_patch_id(ctx);
1392 LLVMValueRef vertices_per_patch = LLVMConstInt(ctx->ac.i32, ctx->tcs_vertices_per_patch, false);
1393 constant16 = LLVMConstInt(ctx->ac.i32, 16, false);
1394 param_stride = calc_param_stride(ctx, vertex_index);
1395 if (vertex_index) {
1396 base_addr = ac_build_imad(&ctx->ac, rel_patch_id,
1397 vertices_per_patch, vertex_index);
1398 } else {
1399 base_addr = rel_patch_id;
1400 }
1401
1402 base_addr = LLVMBuildAdd(ctx->ac.builder, base_addr,
1403 LLVMBuildMul(ctx->ac.builder, param_index,
1404 param_stride, ""), "");
1405
1406 base_addr = LLVMBuildMul(ctx->ac.builder, base_addr, constant16, "");
1407
1408 if (!vertex_index) {
1409 LLVMValueRef patch_data_offset = get_non_vertex_index_offset(ctx);
1410
1411 base_addr = LLVMBuildAdd(ctx->ac.builder, base_addr,
1412 patch_data_offset, "");
1413 }
1414 return base_addr;
1415 }
1416
1417 static LLVMValueRef get_tcs_tes_buffer_address_params(struct radv_shader_context *ctx,
1418 unsigned param,
1419 unsigned const_index,
1420 bool is_compact,
1421 LLVMValueRef vertex_index,
1422 LLVMValueRef indir_index)
1423 {
1424 LLVMValueRef param_index;
1425
1426 if (indir_index)
1427 param_index = LLVMBuildAdd(ctx->ac.builder, LLVMConstInt(ctx->ac.i32, param, false),
1428 indir_index, "");
1429 else {
1430 if (const_index && !is_compact)
1431 param += const_index;
1432 param_index = LLVMConstInt(ctx->ac.i32, param, false);
1433 }
1434 return get_tcs_tes_buffer_address(ctx, vertex_index, param_index);
1435 }
1436
1437 static LLVMValueRef
1438 get_dw_address(struct radv_shader_context *ctx,
1439 LLVMValueRef dw_addr,
1440 unsigned param,
1441 unsigned const_index,
1442 bool compact_const_index,
1443 LLVMValueRef vertex_index,
1444 LLVMValueRef stride,
1445 LLVMValueRef indir_index)
1446
1447 {
1448
1449 if (vertex_index) {
1450 dw_addr = LLVMBuildAdd(ctx->ac.builder, dw_addr,
1451 LLVMBuildMul(ctx->ac.builder,
1452 vertex_index,
1453 stride, ""), "");
1454 }
1455
1456 if (indir_index)
1457 dw_addr = LLVMBuildAdd(ctx->ac.builder, dw_addr,
1458 LLVMBuildMul(ctx->ac.builder, indir_index,
1459 LLVMConstInt(ctx->ac.i32, 4, false), ""), "");
1460 else if (const_index && !compact_const_index)
1461 dw_addr = LLVMBuildAdd(ctx->ac.builder, dw_addr,
1462 LLVMConstInt(ctx->ac.i32, const_index * 4, false), "");
1463
1464 dw_addr = LLVMBuildAdd(ctx->ac.builder, dw_addr,
1465 LLVMConstInt(ctx->ac.i32, param * 4, false), "");
1466
1467 if (const_index && compact_const_index)
1468 dw_addr = LLVMBuildAdd(ctx->ac.builder, dw_addr,
1469 LLVMConstInt(ctx->ac.i32, const_index, false), "");
1470 return dw_addr;
1471 }
1472
1473 static LLVMValueRef
1474 load_tcs_varyings(struct ac_shader_abi *abi,
1475 LLVMTypeRef type,
1476 LLVMValueRef vertex_index,
1477 LLVMValueRef indir_index,
1478 unsigned const_index,
1479 unsigned location,
1480 unsigned driver_location,
1481 unsigned component,
1482 unsigned num_components,
1483 bool is_patch,
1484 bool is_compact,
1485 bool load_input)
1486 {
1487 struct radv_shader_context *ctx = radv_shader_context_from_abi(abi);
1488 LLVMValueRef dw_addr, stride;
1489 LLVMValueRef value[4], result;
1490 unsigned param = shader_io_get_unique_index(location);
1491
1492 if (load_input) {
1493 uint32_t input_vertex_size = (ctx->tcs_num_inputs * 16) / 4;
1494 stride = LLVMConstInt(ctx->ac.i32, input_vertex_size, false);
1495 dw_addr = get_tcs_in_current_patch_offset(ctx);
1496 } else {
1497 if (!is_patch) {
1498 stride = get_tcs_out_vertex_stride(ctx);
1499 dw_addr = get_tcs_out_current_patch_offset(ctx);
1500 } else {
1501 dw_addr = get_tcs_out_current_patch_data_offset(ctx);
1502 stride = NULL;
1503 }
1504 }
1505
1506 dw_addr = get_dw_address(ctx, dw_addr, param, const_index, is_compact, vertex_index, stride,
1507 indir_index);
1508
1509 for (unsigned i = 0; i < num_components + component; i++) {
1510 value[i] = ac_lds_load(&ctx->ac, dw_addr);
1511 dw_addr = LLVMBuildAdd(ctx->ac.builder, dw_addr,
1512 ctx->ac.i32_1, "");
1513 }
1514 result = ac_build_varying_gather_values(&ctx->ac, value, num_components, component);
1515 return result;
1516 }
1517
1518 static void
1519 store_tcs_output(struct ac_shader_abi *abi,
1520 const nir_variable *var,
1521 LLVMValueRef vertex_index,
1522 LLVMValueRef param_index,
1523 unsigned const_index,
1524 LLVMValueRef src,
1525 unsigned writemask)
1526 {
1527 struct radv_shader_context *ctx = radv_shader_context_from_abi(abi);
1528 const unsigned location = var->data.location;
1529 unsigned component = var->data.location_frac;
1530 const bool is_patch = var->data.patch;
1531 const bool is_compact = var->data.compact;
1532 LLVMValueRef dw_addr;
1533 LLVMValueRef stride = NULL;
1534 LLVMValueRef buf_addr = NULL;
1535 unsigned param;
1536 bool store_lds = true;
1537
1538 if (is_patch) {
1539 if (!(ctx->tcs_patch_outputs_read & (1U << (location - VARYING_SLOT_PATCH0))))
1540 store_lds = false;
1541 } else {
1542 if (!(ctx->tcs_outputs_read & (1ULL << location)))
1543 store_lds = false;
1544 }
1545
1546 param = shader_io_get_unique_index(location);
1547 if ((location == VARYING_SLOT_CLIP_DIST0 || location == VARYING_SLOT_CLIP_DIST1) && is_compact) {
1548 const_index += component;
1549 component = 0;
1550
1551 if (const_index >= 4) {
1552 const_index -= 4;
1553 param++;
1554 }
1555 }
1556
1557 if (!is_patch) {
1558 stride = get_tcs_out_vertex_stride(ctx);
1559 dw_addr = get_tcs_out_current_patch_offset(ctx);
1560 } else {
1561 dw_addr = get_tcs_out_current_patch_data_offset(ctx);
1562 }
1563
1564 dw_addr = get_dw_address(ctx, dw_addr, param, const_index, is_compact, vertex_index, stride,
1565 param_index);
1566 buf_addr = get_tcs_tes_buffer_address_params(ctx, param, const_index, is_compact,
1567 vertex_index, param_index);
1568
1569 bool is_tess_factor = false;
1570 if (location == VARYING_SLOT_TESS_LEVEL_INNER ||
1571 location == VARYING_SLOT_TESS_LEVEL_OUTER)
1572 is_tess_factor = true;
1573
1574 unsigned base = is_compact ? const_index : 0;
1575 for (unsigned chan = 0; chan < 8; chan++) {
1576 if (!(writemask & (1 << chan)))
1577 continue;
1578 LLVMValueRef value = ac_llvm_extract_elem(&ctx->ac, src, chan - component);
1579 value = ac_to_integer(&ctx->ac, value);
1580 value = LLVMBuildZExtOrBitCast(ctx->ac.builder, value, ctx->ac.i32, "");
1581
1582 if (store_lds || is_tess_factor) {
1583 LLVMValueRef dw_addr_chan =
1584 LLVMBuildAdd(ctx->ac.builder, dw_addr,
1585 LLVMConstInt(ctx->ac.i32, chan, false), "");
1586 ac_lds_store(&ctx->ac, dw_addr_chan, value);
1587 }
1588
1589 if (!is_tess_factor && writemask != 0xF)
1590 ac_build_buffer_store_dword(&ctx->ac, ctx->hs_ring_tess_offchip, value, 1,
1591 buf_addr, ctx->oc_lds,
1592 4 * (base + chan), 1, 0, true, false);
1593 }
1594
1595 if (writemask == 0xF) {
1596 ac_build_buffer_store_dword(&ctx->ac, ctx->hs_ring_tess_offchip, src, 4,
1597 buf_addr, ctx->oc_lds,
1598 (base * 4), 1, 0, true, false);
1599 }
1600 }
1601
1602 static LLVMValueRef
1603 load_tes_input(struct ac_shader_abi *abi,
1604 LLVMTypeRef type,
1605 LLVMValueRef vertex_index,
1606 LLVMValueRef param_index,
1607 unsigned const_index,
1608 unsigned location,
1609 unsigned driver_location,
1610 unsigned component,
1611 unsigned num_components,
1612 bool is_patch,
1613 bool is_compact,
1614 bool load_input)
1615 {
1616 struct radv_shader_context *ctx = radv_shader_context_from_abi(abi);
1617 LLVMValueRef buf_addr;
1618 LLVMValueRef result;
1619 unsigned param = shader_io_get_unique_index(location);
1620
1621 if ((location == VARYING_SLOT_CLIP_DIST0 || location == VARYING_SLOT_CLIP_DIST1) && is_compact) {
1622 const_index += component;
1623 component = 0;
1624 if (const_index >= 4) {
1625 const_index -= 4;
1626 param++;
1627 }
1628 }
1629
1630 buf_addr = get_tcs_tes_buffer_address_params(ctx, param, const_index,
1631 is_compact, vertex_index, param_index);
1632
1633 LLVMValueRef comp_offset = LLVMConstInt(ctx->ac.i32, component * 4, false);
1634 buf_addr = LLVMBuildAdd(ctx->ac.builder, buf_addr, comp_offset, "");
1635
1636 result = ac_build_buffer_load(&ctx->ac, ctx->hs_ring_tess_offchip, num_components, NULL,
1637 buf_addr, ctx->oc_lds, is_compact ? (4 * const_index) : 0, 1, 0, true, false);
1638 result = ac_trim_vector(&ctx->ac, result, num_components);
1639 return result;
1640 }
1641
1642 static LLVMValueRef
1643 load_gs_input(struct ac_shader_abi *abi,
1644 unsigned location,
1645 unsigned driver_location,
1646 unsigned component,
1647 unsigned num_components,
1648 unsigned vertex_index,
1649 unsigned const_index,
1650 LLVMTypeRef type)
1651 {
1652 struct radv_shader_context *ctx = radv_shader_context_from_abi(abi);
1653 LLVMValueRef vtx_offset;
1654 unsigned param, vtx_offset_param;
1655 LLVMValueRef value[4], result;
1656
1657 vtx_offset_param = vertex_index;
1658 assert(vtx_offset_param < 6);
1659 vtx_offset = LLVMBuildMul(ctx->ac.builder, ctx->gs_vtx_offset[vtx_offset_param],
1660 LLVMConstInt(ctx->ac.i32, 4, false), "");
1661
1662 param = shader_io_get_unique_index(location);
1663
1664 for (unsigned i = component; i < num_components + component; i++) {
1665 if (ctx->ac.chip_class >= GFX9) {
1666 LLVMValueRef dw_addr = ctx->gs_vtx_offset[vtx_offset_param];
1667 dw_addr = LLVMBuildAdd(ctx->ac.builder, dw_addr,
1668 LLVMConstInt(ctx->ac.i32, param * 4 + i + const_index, 0), "");
1669 value[i] = ac_lds_load(&ctx->ac, dw_addr);
1670 } else {
1671 LLVMValueRef soffset =
1672 LLVMConstInt(ctx->ac.i32,
1673 (param * 4 + i + const_index) * 256,
1674 false);
1675
1676 value[i] = ac_build_buffer_load(&ctx->ac,
1677 ctx->esgs_ring, 1,
1678 ctx->ac.i32_0,
1679 vtx_offset, soffset,
1680 0, 1, 0, true, false);
1681 }
1682
1683 if (ac_get_type_size(type) == 2) {
1684 value[i] = LLVMBuildBitCast(ctx->ac.builder, value[i], ctx->ac.i32, "");
1685 value[i] = LLVMBuildTrunc(ctx->ac.builder, value[i], ctx->ac.i16, "");
1686 }
1687 value[i] = LLVMBuildBitCast(ctx->ac.builder, value[i], type, "");
1688 }
1689 result = ac_build_varying_gather_values(&ctx->ac, value, num_components, component);
1690 result = ac_to_integer(&ctx->ac, result);
1691 return result;
1692 }
1693
1694
1695 static void radv_emit_kill(struct ac_shader_abi *abi, LLVMValueRef visible)
1696 {
1697 struct radv_shader_context *ctx = radv_shader_context_from_abi(abi);
1698 ac_build_kill_if_false(&ctx->ac, visible);
1699 }
1700
1701 static LLVMValueRef lookup_interp_param(struct ac_shader_abi *abi,
1702 enum glsl_interp_mode interp, unsigned location)
1703 {
1704 struct radv_shader_context *ctx = radv_shader_context_from_abi(abi);
1705
1706 switch (interp) {
1707 case INTERP_MODE_FLAT:
1708 default:
1709 return NULL;
1710 case INTERP_MODE_SMOOTH:
1711 case INTERP_MODE_NONE:
1712 if (location == INTERP_CENTER)
1713 return ctx->persp_center;
1714 else if (location == INTERP_CENTROID)
1715 return ctx->persp_centroid;
1716 else if (location == INTERP_SAMPLE)
1717 return ctx->persp_sample;
1718 break;
1719 case INTERP_MODE_NOPERSPECTIVE:
1720 if (location == INTERP_CENTER)
1721 return ctx->linear_center;
1722 else if (location == INTERP_CENTROID)
1723 return ctx->linear_centroid;
1724 else if (location == INTERP_SAMPLE)
1725 return ctx->linear_sample;
1726 break;
1727 }
1728 return NULL;
1729 }
1730
1731 static uint32_t
1732 radv_get_sample_pos_offset(uint32_t num_samples)
1733 {
1734 uint32_t sample_pos_offset = 0;
1735
1736 switch (num_samples) {
1737 case 2:
1738 sample_pos_offset = 1;
1739 break;
1740 case 4:
1741 sample_pos_offset = 3;
1742 break;
1743 case 8:
1744 sample_pos_offset = 7;
1745 break;
1746 default:
1747 break;
1748 }
1749 return sample_pos_offset;
1750 }
1751
1752 static LLVMValueRef load_sample_position(struct ac_shader_abi *abi,
1753 LLVMValueRef sample_id)
1754 {
1755 struct radv_shader_context *ctx = radv_shader_context_from_abi(abi);
1756
1757 LLVMValueRef result;
1758 LLVMValueRef ptr = ac_build_gep0(&ctx->ac, ctx->ring_offsets, LLVMConstInt(ctx->ac.i32, RING_PS_SAMPLE_POSITIONS, false));
1759
1760 ptr = LLVMBuildBitCast(ctx->ac.builder, ptr,
1761 ac_array_in_const_addr_space(ctx->ac.v2f32), "");
1762
1763 uint32_t sample_pos_offset =
1764 radv_get_sample_pos_offset(ctx->options->key.fs.num_samples);
1765
1766 sample_id =
1767 LLVMBuildAdd(ctx->ac.builder, sample_id,
1768 LLVMConstInt(ctx->ac.i32, sample_pos_offset, false), "");
1769 result = ac_build_load_invariant(&ctx->ac, ptr, sample_id);
1770
1771 return result;
1772 }
1773
1774
1775 static LLVMValueRef load_sample_mask_in(struct ac_shader_abi *abi)
1776 {
1777 struct radv_shader_context *ctx = radv_shader_context_from_abi(abi);
1778 uint8_t log2_ps_iter_samples;
1779
1780 if (ctx->shader_info->info.ps.force_persample) {
1781 log2_ps_iter_samples =
1782 util_logbase2(ctx->options->key.fs.num_samples);
1783 } else {
1784 log2_ps_iter_samples = ctx->options->key.fs.log2_ps_iter_samples;
1785 }
1786
1787 /* The bit pattern matches that used by fixed function fragment
1788 * processing. */
1789 static const uint16_t ps_iter_masks[] = {
1790 0xffff, /* not used */
1791 0x5555,
1792 0x1111,
1793 0x0101,
1794 0x0001,
1795 };
1796 assert(log2_ps_iter_samples < ARRAY_SIZE(ps_iter_masks));
1797
1798 uint32_t ps_iter_mask = ps_iter_masks[log2_ps_iter_samples];
1799
1800 LLVMValueRef result, sample_id;
1801 sample_id = ac_unpack_param(&ctx->ac, abi->ancillary, 8, 4);
1802 sample_id = LLVMBuildShl(ctx->ac.builder, LLVMConstInt(ctx->ac.i32, ps_iter_mask, false), sample_id, "");
1803 result = LLVMBuildAnd(ctx->ac.builder, sample_id, abi->sample_coverage, "");
1804 return result;
1805 }
1806
1807
1808 static void
1809 visit_emit_vertex(struct ac_shader_abi *abi, unsigned stream, LLVMValueRef *addrs)
1810 {
1811 LLVMValueRef gs_next_vertex;
1812 LLVMValueRef can_emit;
1813 unsigned offset = 0;
1814 struct radv_shader_context *ctx = radv_shader_context_from_abi(abi);
1815
1816 /* Write vertex attribute values to GSVS ring */
1817 gs_next_vertex = LLVMBuildLoad(ctx->ac.builder,
1818 ctx->gs_next_vertex[stream],
1819 "");
1820
1821 /* If this thread has already emitted the declared maximum number of
1822 * vertices, kill it: excessive vertex emissions are not supposed to
1823 * have any effect, and GS threads have no externally observable
1824 * effects other than emitting vertices.
1825 */
1826 can_emit = LLVMBuildICmp(ctx->ac.builder, LLVMIntULT, gs_next_vertex,
1827 LLVMConstInt(ctx->ac.i32, ctx->gs_max_out_vertices, false), "");
1828 ac_build_kill_if_false(&ctx->ac, can_emit);
1829
1830 for (unsigned i = 0; i < AC_LLVM_MAX_OUTPUTS; ++i) {
1831 unsigned output_usage_mask =
1832 ctx->shader_info->info.gs.output_usage_mask[i];
1833 uint8_t output_stream =
1834 ctx->shader_info->info.gs.output_streams[i];
1835 LLVMValueRef *out_ptr = &addrs[i * 4];
1836 int length = util_last_bit(output_usage_mask);
1837
1838 if (!(ctx->output_mask & (1ull << i)) ||
1839 output_stream != stream)
1840 continue;
1841
1842 for (unsigned j = 0; j < length; j++) {
1843 if (!(output_usage_mask & (1 << j)))
1844 continue;
1845
1846 LLVMValueRef out_val = LLVMBuildLoad(ctx->ac.builder,
1847 out_ptr[j], "");
1848 LLVMValueRef voffset =
1849 LLVMConstInt(ctx->ac.i32, offset *
1850 ctx->gs_max_out_vertices, false);
1851
1852 offset++;
1853
1854 voffset = LLVMBuildAdd(ctx->ac.builder, voffset, gs_next_vertex, "");
1855 voffset = LLVMBuildMul(ctx->ac.builder, voffset, LLVMConstInt(ctx->ac.i32, 4, false), "");
1856
1857 out_val = ac_to_integer(&ctx->ac, out_val);
1858 out_val = LLVMBuildZExtOrBitCast(ctx->ac.builder, out_val, ctx->ac.i32, "");
1859
1860 ac_build_buffer_store_dword(&ctx->ac,
1861 ctx->gsvs_ring[stream],
1862 out_val, 1,
1863 voffset, ctx->gs2vs_offset, 0,
1864 1, 1, true, true);
1865 }
1866 }
1867
1868 gs_next_vertex = LLVMBuildAdd(ctx->ac.builder, gs_next_vertex,
1869 ctx->ac.i32_1, "");
1870 LLVMBuildStore(ctx->ac.builder, gs_next_vertex, ctx->gs_next_vertex[stream]);
1871
1872 ac_build_sendmsg(&ctx->ac,
1873 AC_SENDMSG_GS_OP_EMIT | AC_SENDMSG_GS | (stream << 8),
1874 ctx->gs_wave_id);
1875 }
1876
1877 static void
1878 visit_end_primitive(struct ac_shader_abi *abi, unsigned stream)
1879 {
1880 struct radv_shader_context *ctx = radv_shader_context_from_abi(abi);
1881 ac_build_sendmsg(&ctx->ac, AC_SENDMSG_GS_OP_CUT | AC_SENDMSG_GS | (stream << 8), ctx->gs_wave_id);
1882 }
1883
1884 static LLVMValueRef
1885 load_tess_coord(struct ac_shader_abi *abi)
1886 {
1887 struct radv_shader_context *ctx = radv_shader_context_from_abi(abi);
1888
1889 LLVMValueRef coord[4] = {
1890 ctx->tes_u,
1891 ctx->tes_v,
1892 ctx->ac.f32_0,
1893 ctx->ac.f32_0,
1894 };
1895
1896 if (ctx->tes_primitive_mode == GL_TRIANGLES)
1897 coord[2] = LLVMBuildFSub(ctx->ac.builder, ctx->ac.f32_1,
1898 LLVMBuildFAdd(ctx->ac.builder, coord[0], coord[1], ""), "");
1899
1900 return ac_build_gather_values(&ctx->ac, coord, 3);
1901 }
1902
1903 static LLVMValueRef
1904 load_patch_vertices_in(struct ac_shader_abi *abi)
1905 {
1906 struct radv_shader_context *ctx = radv_shader_context_from_abi(abi);
1907 return LLVMConstInt(ctx->ac.i32, ctx->options->key.tcs.input_vertices, false);
1908 }
1909
1910
1911 static LLVMValueRef radv_load_base_vertex(struct ac_shader_abi *abi)
1912 {
1913 return abi->base_vertex;
1914 }
1915
1916 static LLVMValueRef radv_load_ssbo(struct ac_shader_abi *abi,
1917 LLVMValueRef buffer_ptr, bool write)
1918 {
1919 struct radv_shader_context *ctx = radv_shader_context_from_abi(abi);
1920 LLVMValueRef result;
1921
1922 LLVMSetMetadata(buffer_ptr, ctx->ac.uniform_md_kind, ctx->ac.empty_md);
1923
1924 result = LLVMBuildLoad(ctx->ac.builder, buffer_ptr, "");
1925 LLVMSetMetadata(result, ctx->ac.invariant_load_md_kind, ctx->ac.empty_md);
1926
1927 return result;
1928 }
1929
1930 static LLVMValueRef radv_load_ubo(struct ac_shader_abi *abi, LLVMValueRef buffer_ptr)
1931 {
1932 struct radv_shader_context *ctx = radv_shader_context_from_abi(abi);
1933 LLVMValueRef result;
1934
1935 if (LLVMGetTypeKind(LLVMTypeOf(buffer_ptr)) != LLVMPointerTypeKind) {
1936 /* Do not load the descriptor for inlined uniform blocks. */
1937 return buffer_ptr;
1938 }
1939
1940 LLVMSetMetadata(buffer_ptr, ctx->ac.uniform_md_kind, ctx->ac.empty_md);
1941
1942 result = LLVMBuildLoad(ctx->ac.builder, buffer_ptr, "");
1943 LLVMSetMetadata(result, ctx->ac.invariant_load_md_kind, ctx->ac.empty_md);
1944
1945 return result;
1946 }
1947
1948 static LLVMValueRef radv_get_sampler_desc(struct ac_shader_abi *abi,
1949 unsigned descriptor_set,
1950 unsigned base_index,
1951 unsigned constant_index,
1952 LLVMValueRef index,
1953 enum ac_descriptor_type desc_type,
1954 bool image, bool write,
1955 bool bindless)
1956 {
1957 struct radv_shader_context *ctx = radv_shader_context_from_abi(abi);
1958 LLVMValueRef list = ctx->descriptor_sets[descriptor_set];
1959 struct radv_descriptor_set_layout *layout = ctx->options->layout->set[descriptor_set].layout;
1960 struct radv_descriptor_set_binding_layout *binding = layout->binding + base_index;
1961 unsigned offset = binding->offset;
1962 unsigned stride = binding->size;
1963 unsigned type_size;
1964 LLVMBuilderRef builder = ctx->ac.builder;
1965 LLVMTypeRef type;
1966
1967 assert(base_index < layout->binding_count);
1968
1969 switch (desc_type) {
1970 case AC_DESC_IMAGE:
1971 type = ctx->ac.v8i32;
1972 type_size = 32;
1973 break;
1974 case AC_DESC_FMASK:
1975 type = ctx->ac.v8i32;
1976 offset += 32;
1977 type_size = 32;
1978 break;
1979 case AC_DESC_SAMPLER:
1980 type = ctx->ac.v4i32;
1981 if (binding->type == VK_DESCRIPTOR_TYPE_COMBINED_IMAGE_SAMPLER) {
1982 offset += radv_combined_image_descriptor_sampler_offset(binding);
1983 }
1984
1985 type_size = 16;
1986 break;
1987 case AC_DESC_BUFFER:
1988 type = ctx->ac.v4i32;
1989 type_size = 16;
1990 break;
1991 case AC_DESC_PLANE_0:
1992 case AC_DESC_PLANE_1:
1993 case AC_DESC_PLANE_2:
1994 type = ctx->ac.v8i32;
1995 type_size = 32;
1996 offset += 32 * (desc_type - AC_DESC_PLANE_0);
1997 break;
1998 default:
1999 unreachable("invalid desc_type\n");
2000 }
2001
2002 offset += constant_index * stride;
2003
2004 if (desc_type == AC_DESC_SAMPLER && binding->immutable_samplers_offset &&
2005 (!index || binding->immutable_samplers_equal)) {
2006 if (binding->immutable_samplers_equal)
2007 constant_index = 0;
2008
2009 const uint32_t *samplers = radv_immutable_samplers(layout, binding);
2010
2011 LLVMValueRef constants[] = {
2012 LLVMConstInt(ctx->ac.i32, samplers[constant_index * 4 + 0], 0),
2013 LLVMConstInt(ctx->ac.i32, samplers[constant_index * 4 + 1], 0),
2014 LLVMConstInt(ctx->ac.i32, samplers[constant_index * 4 + 2], 0),
2015 LLVMConstInt(ctx->ac.i32, samplers[constant_index * 4 + 3], 0),
2016 };
2017 return ac_build_gather_values(&ctx->ac, constants, 4);
2018 }
2019
2020 assert(stride % type_size == 0);
2021
2022 if (!index)
2023 index = ctx->ac.i32_0;
2024
2025 index = LLVMBuildMul(builder, index, LLVMConstInt(ctx->ac.i32, stride / type_size, 0), "");
2026
2027 list = ac_build_gep0(&ctx->ac, list, LLVMConstInt(ctx->ac.i32, offset, 0));
2028 list = LLVMBuildPointerCast(builder, list,
2029 ac_array_in_const32_addr_space(type), "");
2030
2031 return ac_build_load_to_sgpr(&ctx->ac, list, index);
2032 }
2033
2034 /* For 2_10_10_10 formats the alpha is handled as unsigned by pre-vega HW.
2035 * so we may need to fix it up. */
2036 static LLVMValueRef
2037 adjust_vertex_fetch_alpha(struct radv_shader_context *ctx,
2038 unsigned adjustment,
2039 LLVMValueRef alpha)
2040 {
2041 if (adjustment == RADV_ALPHA_ADJUST_NONE)
2042 return alpha;
2043
2044 LLVMValueRef c30 = LLVMConstInt(ctx->ac.i32, 30, 0);
2045
2046 alpha = LLVMBuildBitCast(ctx->ac.builder, alpha, ctx->ac.f32, "");
2047
2048 if (adjustment == RADV_ALPHA_ADJUST_SSCALED)
2049 alpha = LLVMBuildFPToUI(ctx->ac.builder, alpha, ctx->ac.i32, "");
2050 else
2051 alpha = ac_to_integer(&ctx->ac, alpha);
2052
2053 /* For the integer-like cases, do a natural sign extension.
2054 *
2055 * For the SNORM case, the values are 0.0, 0.333, 0.666, 1.0
2056 * and happen to contain 0, 1, 2, 3 as the two LSBs of the
2057 * exponent.
2058 */
2059 alpha = LLVMBuildShl(ctx->ac.builder, alpha,
2060 adjustment == RADV_ALPHA_ADJUST_SNORM ?
2061 LLVMConstInt(ctx->ac.i32, 7, 0) : c30, "");
2062 alpha = LLVMBuildAShr(ctx->ac.builder, alpha, c30, "");
2063
2064 /* Convert back to the right type. */
2065 if (adjustment == RADV_ALPHA_ADJUST_SNORM) {
2066 LLVMValueRef clamp;
2067 LLVMValueRef neg_one = LLVMConstReal(ctx->ac.f32, -1.0);
2068 alpha = LLVMBuildSIToFP(ctx->ac.builder, alpha, ctx->ac.f32, "");
2069 clamp = LLVMBuildFCmp(ctx->ac.builder, LLVMRealULT, alpha, neg_one, "");
2070 alpha = LLVMBuildSelect(ctx->ac.builder, clamp, neg_one, alpha, "");
2071 } else if (adjustment == RADV_ALPHA_ADJUST_SSCALED) {
2072 alpha = LLVMBuildSIToFP(ctx->ac.builder, alpha, ctx->ac.f32, "");
2073 }
2074
2075 return LLVMBuildBitCast(ctx->ac.builder, alpha, ctx->ac.i32, "");
2076 }
2077
2078 static unsigned
2079 get_num_channels_from_data_format(unsigned data_format)
2080 {
2081 switch (data_format) {
2082 case V_008F0C_BUF_DATA_FORMAT_8:
2083 case V_008F0C_BUF_DATA_FORMAT_16:
2084 case V_008F0C_BUF_DATA_FORMAT_32:
2085 return 1;
2086 case V_008F0C_BUF_DATA_FORMAT_8_8:
2087 case V_008F0C_BUF_DATA_FORMAT_16_16:
2088 case V_008F0C_BUF_DATA_FORMAT_32_32:
2089 return 2;
2090 case V_008F0C_BUF_DATA_FORMAT_10_11_11:
2091 case V_008F0C_BUF_DATA_FORMAT_11_11_10:
2092 case V_008F0C_BUF_DATA_FORMAT_32_32_32:
2093 return 3;
2094 case V_008F0C_BUF_DATA_FORMAT_8_8_8_8:
2095 case V_008F0C_BUF_DATA_FORMAT_10_10_10_2:
2096 case V_008F0C_BUF_DATA_FORMAT_2_10_10_10:
2097 case V_008F0C_BUF_DATA_FORMAT_16_16_16_16:
2098 case V_008F0C_BUF_DATA_FORMAT_32_32_32_32:
2099 return 4;
2100 default:
2101 break;
2102 }
2103
2104 return 4;
2105 }
2106
2107 static LLVMValueRef
2108 radv_fixup_vertex_input_fetches(struct radv_shader_context *ctx,
2109 LLVMValueRef value,
2110 unsigned num_channels,
2111 bool is_float)
2112 {
2113 LLVMValueRef zero = is_float ? ctx->ac.f32_0 : ctx->ac.i32_0;
2114 LLVMValueRef one = is_float ? ctx->ac.f32_1 : ctx->ac.i32_1;
2115 LLVMValueRef chan[4];
2116
2117 if (LLVMGetTypeKind(LLVMTypeOf(value)) == LLVMVectorTypeKind) {
2118 unsigned vec_size = LLVMGetVectorSize(LLVMTypeOf(value));
2119
2120 if (num_channels == 4 && num_channels == vec_size)
2121 return value;
2122
2123 num_channels = MIN2(num_channels, vec_size);
2124
2125 for (unsigned i = 0; i < num_channels; i++)
2126 chan[i] = ac_llvm_extract_elem(&ctx->ac, value, i);
2127 } else {
2128 if (num_channels) {
2129 assert(num_channels == 1);
2130 chan[0] = value;
2131 }
2132 }
2133
2134 for (unsigned i = num_channels; i < 4; i++) {
2135 chan[i] = i == 3 ? one : zero;
2136 chan[i] = ac_to_integer(&ctx->ac, chan[i]);
2137 }
2138
2139 return ac_build_gather_values(&ctx->ac, chan, 4);
2140 }
2141
2142 static void
2143 handle_vs_input_decl(struct radv_shader_context *ctx,
2144 struct nir_variable *variable)
2145 {
2146 LLVMValueRef t_list_ptr = ctx->vertex_buffers;
2147 LLVMValueRef t_offset;
2148 LLVMValueRef t_list;
2149 LLVMValueRef input;
2150 LLVMValueRef buffer_index;
2151 unsigned attrib_count = glsl_count_attribute_slots(variable->type, true);
2152 uint8_t input_usage_mask =
2153 ctx->shader_info->info.vs.input_usage_mask[variable->data.location];
2154 unsigned num_input_channels = util_last_bit(input_usage_mask);
2155
2156 variable->data.driver_location = variable->data.location * 4;
2157
2158 enum glsl_base_type type = glsl_get_base_type(variable->type);
2159 for (unsigned i = 0; i < attrib_count; ++i) {
2160 LLVMValueRef output[4];
2161 unsigned attrib_index = variable->data.location + i - VERT_ATTRIB_GENERIC0;
2162 unsigned attrib_format = ctx->options->key.vs.vertex_attribute_formats[attrib_index];
2163 unsigned data_format = attrib_format & 0x0f;
2164 unsigned num_format = (attrib_format >> 4) & 0x07;
2165 bool is_float = num_format != V_008F0C_BUF_NUM_FORMAT_UINT &&
2166 num_format != V_008F0C_BUF_NUM_FORMAT_SINT;
2167
2168 if (ctx->options->key.vs.instance_rate_inputs & (1u << attrib_index)) {
2169 uint32_t divisor = ctx->options->key.vs.instance_rate_divisors[attrib_index];
2170
2171 if (divisor) {
2172 buffer_index = ctx->abi.instance_id;
2173
2174 if (divisor != 1) {
2175 buffer_index = LLVMBuildUDiv(ctx->ac.builder, buffer_index,
2176 LLVMConstInt(ctx->ac.i32, divisor, 0), "");
2177 }
2178
2179 if (ctx->options->key.vs.as_ls) {
2180 ctx->shader_info->vs.vgpr_comp_cnt =
2181 MAX2(2, ctx->shader_info->vs.vgpr_comp_cnt);
2182 } else {
2183 ctx->shader_info->vs.vgpr_comp_cnt =
2184 MAX2(1, ctx->shader_info->vs.vgpr_comp_cnt);
2185 }
2186 } else {
2187 buffer_index = ctx->ac.i32_0;
2188 }
2189
2190 buffer_index = LLVMBuildAdd(ctx->ac.builder, ctx->abi.start_instance, buffer_index, "");
2191 } else
2192 buffer_index = LLVMBuildAdd(ctx->ac.builder, ctx->abi.vertex_id,
2193 ctx->abi.base_vertex, "");
2194
2195 /* Adjust the number of channels to load based on the vertex
2196 * attribute format.
2197 */
2198 unsigned num_format_channels = get_num_channels_from_data_format(data_format);
2199 unsigned num_channels = MIN2(num_input_channels, num_format_channels);
2200 unsigned attrib_binding = ctx->options->key.vs.vertex_attribute_bindings[attrib_index];
2201 unsigned attrib_offset = ctx->options->key.vs.vertex_attribute_offsets[attrib_index];
2202 unsigned attrib_stride = ctx->options->key.vs.vertex_attribute_strides[attrib_index];
2203
2204 if (ctx->options->key.vs.post_shuffle & (1 << attrib_index)) {
2205 /* Always load, at least, 3 channels for formats that
2206 * need to be shuffled because X<->Z.
2207 */
2208 num_channels = MAX2(num_channels, 3);
2209 }
2210
2211 if (attrib_stride != 0 && attrib_offset > attrib_stride) {
2212 LLVMValueRef buffer_offset =
2213 LLVMConstInt(ctx->ac.i32,
2214 attrib_offset / attrib_stride, false);
2215
2216 buffer_index = LLVMBuildAdd(ctx->ac.builder,
2217 buffer_index,
2218 buffer_offset, "");
2219
2220 attrib_offset = attrib_offset % attrib_stride;
2221 }
2222
2223 t_offset = LLVMConstInt(ctx->ac.i32, attrib_binding, false);
2224 t_list = ac_build_load_to_sgpr(&ctx->ac, t_list_ptr, t_offset);
2225
2226 input = ac_build_struct_tbuffer_load(&ctx->ac, t_list,
2227 buffer_index,
2228 LLVMConstInt(ctx->ac.i32, attrib_offset, false),
2229 ctx->ac.i32_0, ctx->ac.i32_0,
2230 num_channels,
2231 data_format, num_format,
2232 false, false, true);
2233
2234 if (ctx->options->key.vs.post_shuffle & (1 << attrib_index)) {
2235 LLVMValueRef c[4];
2236 c[0] = ac_llvm_extract_elem(&ctx->ac, input, 2);
2237 c[1] = ac_llvm_extract_elem(&ctx->ac, input, 1);
2238 c[2] = ac_llvm_extract_elem(&ctx->ac, input, 0);
2239 c[3] = ac_llvm_extract_elem(&ctx->ac, input, 3);
2240
2241 input = ac_build_gather_values(&ctx->ac, c, 4);
2242 }
2243
2244 input = radv_fixup_vertex_input_fetches(ctx, input, num_channels,
2245 is_float);
2246
2247 for (unsigned chan = 0; chan < 4; chan++) {
2248 LLVMValueRef llvm_chan = LLVMConstInt(ctx->ac.i32, chan, false);
2249 output[chan] = LLVMBuildExtractElement(ctx->ac.builder, input, llvm_chan, "");
2250 if (type == GLSL_TYPE_FLOAT16) {
2251 output[chan] = LLVMBuildBitCast(ctx->ac.builder, output[chan], ctx->ac.f32, "");
2252 output[chan] = LLVMBuildFPTrunc(ctx->ac.builder, output[chan], ctx->ac.f16, "");
2253 }
2254 }
2255
2256 unsigned alpha_adjust = (ctx->options->key.vs.alpha_adjust >> (attrib_index * 2)) & 3;
2257 output[3] = adjust_vertex_fetch_alpha(ctx, alpha_adjust, output[3]);
2258
2259 for (unsigned chan = 0; chan < 4; chan++) {
2260 output[chan] = ac_to_integer(&ctx->ac, output[chan]);
2261 if (type == GLSL_TYPE_UINT16 || type == GLSL_TYPE_INT16)
2262 output[chan] = LLVMBuildTrunc(ctx->ac.builder, output[chan], ctx->ac.i16, "");
2263
2264 ctx->inputs[ac_llvm_reg_index_soa(variable->data.location + i, chan)] = output[chan];
2265 }
2266 }
2267 }
2268
2269 static void interp_fs_input(struct radv_shader_context *ctx,
2270 unsigned attr,
2271 LLVMValueRef interp_param,
2272 LLVMValueRef prim_mask,
2273 bool float16,
2274 LLVMValueRef result[4])
2275 {
2276 LLVMValueRef attr_number;
2277 unsigned chan;
2278 LLVMValueRef i, j;
2279 bool interp = !LLVMIsUndef(interp_param);
2280
2281 attr_number = LLVMConstInt(ctx->ac.i32, attr, false);
2282
2283 /* fs.constant returns the param from the middle vertex, so it's not
2284 * really useful for flat shading. It's meant to be used for custom
2285 * interpolation (but the intrinsic can't fetch from the other two
2286 * vertices).
2287 *
2288 * Luckily, it doesn't matter, because we rely on the FLAT_SHADE state
2289 * to do the right thing. The only reason we use fs.constant is that
2290 * fs.interp cannot be used on integers, because they can be equal
2291 * to NaN.
2292 */
2293 if (interp) {
2294 interp_param = LLVMBuildBitCast(ctx->ac.builder, interp_param,
2295 ctx->ac.v2f32, "");
2296
2297 i = LLVMBuildExtractElement(ctx->ac.builder, interp_param,
2298 ctx->ac.i32_0, "");
2299 j = LLVMBuildExtractElement(ctx->ac.builder, interp_param,
2300 ctx->ac.i32_1, "");
2301 }
2302
2303 for (chan = 0; chan < 4; chan++) {
2304 LLVMValueRef llvm_chan = LLVMConstInt(ctx->ac.i32, chan, false);
2305
2306 if (interp && float16) {
2307 result[chan] = ac_build_fs_interp_f16(&ctx->ac,
2308 llvm_chan,
2309 attr_number,
2310 prim_mask, i, j);
2311 } else if (interp) {
2312 result[chan] = ac_build_fs_interp(&ctx->ac,
2313 llvm_chan,
2314 attr_number,
2315 prim_mask, i, j);
2316 } else {
2317 result[chan] = ac_build_fs_interp_mov(&ctx->ac,
2318 LLVMConstInt(ctx->ac.i32, 2, false),
2319 llvm_chan,
2320 attr_number,
2321 prim_mask);
2322 result[chan] = LLVMBuildBitCast(ctx->ac.builder, result[chan], ctx->ac.i32, "");
2323 result[chan] = LLVMBuildTruncOrBitCast(ctx->ac.builder, result[chan], float16 ? ctx->ac.i16 : ctx->ac.i32, "");
2324 }
2325 }
2326 }
2327
2328 static void mark_16bit_fs_input(struct radv_shader_context *ctx,
2329 const struct glsl_type *type,
2330 int location)
2331 {
2332 if (glsl_type_is_scalar(type) || glsl_type_is_vector(type) || glsl_type_is_matrix(type)) {
2333 unsigned attrib_count = glsl_count_attribute_slots(type, false);
2334 if (glsl_type_is_16bit(type)) {
2335 ctx->float16_shaded_mask |= ((1ull << attrib_count) - 1) << location;
2336 }
2337 } else if (glsl_type_is_array(type)) {
2338 unsigned stride = glsl_count_attribute_slots(glsl_get_array_element(type), false);
2339 for (unsigned i = 0; i < glsl_get_length(type); ++i) {
2340 mark_16bit_fs_input(ctx, glsl_get_array_element(type), location + i * stride);
2341 }
2342 } else {
2343 assert(glsl_type_is_struct_or_ifc(type));
2344 for (unsigned i = 0; i < glsl_get_length(type); i++) {
2345 mark_16bit_fs_input(ctx, glsl_get_struct_field(type, i), location);
2346 location += glsl_count_attribute_slots(glsl_get_struct_field(type, i), false);
2347 }
2348 }
2349 }
2350
2351 static void
2352 handle_fs_input_decl(struct radv_shader_context *ctx,
2353 struct nir_variable *variable)
2354 {
2355 int idx = variable->data.location;
2356 unsigned attrib_count = glsl_count_attribute_slots(variable->type, false);
2357 LLVMValueRef interp = NULL;
2358 uint64_t mask;
2359
2360 variable->data.driver_location = idx * 4;
2361
2362
2363 if (variable->data.compact) {
2364 unsigned component_count = variable->data.location_frac +
2365 glsl_get_length(variable->type);
2366 attrib_count = (component_count + 3) / 4;
2367 } else
2368 mark_16bit_fs_input(ctx, variable->type, idx);
2369
2370 mask = ((1ull << attrib_count) - 1) << variable->data.location;
2371
2372 if (glsl_get_base_type(glsl_without_array(variable->type)) == GLSL_TYPE_FLOAT ||
2373 glsl_get_base_type(glsl_without_array(variable->type)) == GLSL_TYPE_FLOAT16 ||
2374 glsl_get_base_type(glsl_without_array(variable->type)) == GLSL_TYPE_STRUCT) {
2375 unsigned interp_type;
2376 if (variable->data.sample)
2377 interp_type = INTERP_SAMPLE;
2378 else if (variable->data.centroid)
2379 interp_type = INTERP_CENTROID;
2380 else
2381 interp_type = INTERP_CENTER;
2382
2383 interp = lookup_interp_param(&ctx->abi, variable->data.interpolation, interp_type);
2384 }
2385 if (interp == NULL)
2386 interp = LLVMGetUndef(ctx->ac.i32);
2387
2388 for (unsigned i = 0; i < attrib_count; ++i)
2389 ctx->inputs[ac_llvm_reg_index_soa(idx + i, 0)] = interp;
2390
2391 ctx->input_mask |= mask;
2392 }
2393
2394 static void
2395 handle_vs_inputs(struct radv_shader_context *ctx,
2396 struct nir_shader *nir) {
2397 nir_foreach_variable(variable, &nir->inputs)
2398 handle_vs_input_decl(ctx, variable);
2399 }
2400
2401 static void
2402 prepare_interp_optimize(struct radv_shader_context *ctx,
2403 struct nir_shader *nir)
2404 {
2405 bool uses_center = false;
2406 bool uses_centroid = false;
2407 nir_foreach_variable(variable, &nir->inputs) {
2408 if (glsl_get_base_type(glsl_without_array(variable->type)) != GLSL_TYPE_FLOAT ||
2409 variable->data.sample)
2410 continue;
2411
2412 if (variable->data.centroid)
2413 uses_centroid = true;
2414 else
2415 uses_center = true;
2416 }
2417
2418 if (uses_center && uses_centroid) {
2419 LLVMValueRef sel = LLVMBuildICmp(ctx->ac.builder, LLVMIntSLT, ctx->abi.prim_mask, ctx->ac.i32_0, "");
2420 ctx->persp_centroid = LLVMBuildSelect(ctx->ac.builder, sel, ctx->persp_center, ctx->persp_centroid, "");
2421 ctx->linear_centroid = LLVMBuildSelect(ctx->ac.builder, sel, ctx->linear_center, ctx->linear_centroid, "");
2422 }
2423 }
2424
2425 static void
2426 handle_fs_inputs(struct radv_shader_context *ctx,
2427 struct nir_shader *nir)
2428 {
2429 prepare_interp_optimize(ctx, nir);
2430
2431 nir_foreach_variable(variable, &nir->inputs)
2432 handle_fs_input_decl(ctx, variable);
2433
2434 unsigned index = 0;
2435
2436 if (ctx->shader_info->info.ps.uses_input_attachments ||
2437 ctx->shader_info->info.needs_multiview_view_index) {
2438 ctx->input_mask |= 1ull << VARYING_SLOT_LAYER;
2439 ctx->inputs[ac_llvm_reg_index_soa(VARYING_SLOT_LAYER, 0)] = LLVMGetUndef(ctx->ac.i32);
2440 }
2441
2442 for (unsigned i = 0; i < RADEON_LLVM_MAX_INPUTS; ++i) {
2443 LLVMValueRef interp_param;
2444 LLVMValueRef *inputs = ctx->inputs +ac_llvm_reg_index_soa(i, 0);
2445
2446 if (!(ctx->input_mask & (1ull << i)))
2447 continue;
2448
2449 if (i >= VARYING_SLOT_VAR0 || i == VARYING_SLOT_PNTC ||
2450 i == VARYING_SLOT_PRIMITIVE_ID || i == VARYING_SLOT_LAYER) {
2451 interp_param = *inputs;
2452 bool float16 = (ctx->float16_shaded_mask >> i) & 1;
2453 interp_fs_input(ctx, index, interp_param, ctx->abi.prim_mask, float16,
2454 inputs);
2455
2456 if (LLVMIsUndef(interp_param))
2457 ctx->shader_info->fs.flat_shaded_mask |= 1u << index;
2458 if (float16)
2459 ctx->shader_info->fs.float16_shaded_mask |= 1u << index;
2460 if (i >= VARYING_SLOT_VAR0)
2461 ctx->abi.fs_input_attr_indices[i - VARYING_SLOT_VAR0] = index;
2462 ++index;
2463 } else if (i == VARYING_SLOT_CLIP_DIST0) {
2464 int length = ctx->shader_info->info.ps.num_input_clips_culls;
2465
2466 for (unsigned j = 0; j < length; j += 4) {
2467 inputs = ctx->inputs + ac_llvm_reg_index_soa(i, j);
2468
2469 interp_param = *inputs;
2470 interp_fs_input(ctx, index, interp_param,
2471 ctx->abi.prim_mask, false, inputs);
2472 ++index;
2473 }
2474 } else if (i == VARYING_SLOT_POS) {
2475 for(int i = 0; i < 3; ++i)
2476 inputs[i] = ctx->abi.frag_pos[i];
2477
2478 inputs[3] = ac_build_fdiv(&ctx->ac, ctx->ac.f32_1,
2479 ctx->abi.frag_pos[3]);
2480 }
2481 }
2482 ctx->shader_info->fs.num_interp = index;
2483 ctx->shader_info->fs.input_mask = ctx->input_mask >> VARYING_SLOT_VAR0;
2484
2485 if (ctx->shader_info->info.needs_multiview_view_index)
2486 ctx->abi.view_index = ctx->inputs[ac_llvm_reg_index_soa(VARYING_SLOT_LAYER, 0)];
2487 }
2488
2489 static void
2490 scan_shader_output_decl(struct radv_shader_context *ctx,
2491 struct nir_variable *variable,
2492 struct nir_shader *shader,
2493 gl_shader_stage stage)
2494 {
2495 int idx = variable->data.location + variable->data.index;
2496 unsigned attrib_count = glsl_count_attribute_slots(variable->type, false);
2497 uint64_t mask_attribs;
2498
2499 variable->data.driver_location = idx * 4;
2500
2501 /* tess ctrl has it's own load/store paths for outputs */
2502 if (stage == MESA_SHADER_TESS_CTRL)
2503 return;
2504
2505 if (variable->data.compact) {
2506 unsigned component_count = variable->data.location_frac +
2507 glsl_get_length(variable->type);
2508 attrib_count = (component_count + 3) / 4;
2509 }
2510
2511 mask_attribs = ((1ull << attrib_count) - 1) << idx;
2512 if (stage == MESA_SHADER_VERTEX ||
2513 stage == MESA_SHADER_TESS_EVAL ||
2514 stage == MESA_SHADER_GEOMETRY) {
2515 if (idx == VARYING_SLOT_CLIP_DIST0) {
2516 if (stage == MESA_SHADER_VERTEX) {
2517 ctx->shader_info->vs.outinfo.clip_dist_mask = (1 << shader->info.clip_distance_array_size) - 1;
2518 ctx->shader_info->vs.outinfo.cull_dist_mask = (1 << shader->info.cull_distance_array_size) - 1;
2519 ctx->shader_info->vs.outinfo.cull_dist_mask <<= shader->info.clip_distance_array_size;
2520 }
2521 if (stage == MESA_SHADER_TESS_EVAL) {
2522 ctx->shader_info->tes.outinfo.clip_dist_mask = (1 << shader->info.clip_distance_array_size) - 1;
2523 ctx->shader_info->tes.outinfo.cull_dist_mask = (1 << shader->info.cull_distance_array_size) - 1;
2524 ctx->shader_info->tes.outinfo.cull_dist_mask <<= shader->info.clip_distance_array_size;
2525 }
2526 }
2527 }
2528
2529 ctx->output_mask |= mask_attribs;
2530 }
2531
2532
2533 /* Initialize arguments for the shader export intrinsic */
2534 static void
2535 si_llvm_init_export_args(struct radv_shader_context *ctx,
2536 LLVMValueRef *values,
2537 unsigned enabled_channels,
2538 unsigned target,
2539 struct ac_export_args *args)
2540 {
2541 /* Specify the channels that are enabled. */
2542 args->enabled_channels = enabled_channels;
2543
2544 /* Specify whether the EXEC mask represents the valid mask */
2545 args->valid_mask = 0;
2546
2547 /* Specify whether this is the last export */
2548 args->done = 0;
2549
2550 /* Specify the target we are exporting */
2551 args->target = target;
2552
2553 args->compr = false;
2554 args->out[0] = LLVMGetUndef(ctx->ac.f32);
2555 args->out[1] = LLVMGetUndef(ctx->ac.f32);
2556 args->out[2] = LLVMGetUndef(ctx->ac.f32);
2557 args->out[3] = LLVMGetUndef(ctx->ac.f32);
2558
2559 if (!values)
2560 return;
2561
2562 bool is_16bit = ac_get_type_size(LLVMTypeOf(values[0])) == 2;
2563 if (ctx->stage == MESA_SHADER_FRAGMENT) {
2564 unsigned index = target - V_008DFC_SQ_EXP_MRT;
2565 unsigned col_format = (ctx->options->key.fs.col_format >> (4 * index)) & 0xf;
2566 bool is_int8 = (ctx->options->key.fs.is_int8 >> index) & 1;
2567 bool is_int10 = (ctx->options->key.fs.is_int10 >> index) & 1;
2568 unsigned chan;
2569
2570 LLVMValueRef (*packf)(struct ac_llvm_context *ctx, LLVMValueRef args[2]) = NULL;
2571 LLVMValueRef (*packi)(struct ac_llvm_context *ctx, LLVMValueRef args[2],
2572 unsigned bits, bool hi) = NULL;
2573
2574 switch(col_format) {
2575 case V_028714_SPI_SHADER_ZERO:
2576 args->enabled_channels = 0; /* writemask */
2577 args->target = V_008DFC_SQ_EXP_NULL;
2578 break;
2579
2580 case V_028714_SPI_SHADER_32_R:
2581 args->enabled_channels = 1;
2582 args->out[0] = values[0];
2583 break;
2584
2585 case V_028714_SPI_SHADER_32_GR:
2586 args->enabled_channels = 0x3;
2587 args->out[0] = values[0];
2588 args->out[1] = values[1];
2589 break;
2590
2591 case V_028714_SPI_SHADER_32_AR:
2592 args->enabled_channels = 0x9;
2593 args->out[0] = values[0];
2594 args->out[3] = values[3];
2595 break;
2596
2597 case V_028714_SPI_SHADER_FP16_ABGR:
2598 args->enabled_channels = 0x5;
2599 packf = ac_build_cvt_pkrtz_f16;
2600 if (is_16bit) {
2601 for (unsigned chan = 0; chan < 4; chan++)
2602 values[chan] = LLVMBuildFPExt(ctx->ac.builder,
2603 values[chan],
2604 ctx->ac.f32, "");
2605 }
2606 break;
2607
2608 case V_028714_SPI_SHADER_UNORM16_ABGR:
2609 args->enabled_channels = 0x5;
2610 packf = ac_build_cvt_pknorm_u16;
2611 break;
2612
2613 case V_028714_SPI_SHADER_SNORM16_ABGR:
2614 args->enabled_channels = 0x5;
2615 packf = ac_build_cvt_pknorm_i16;
2616 break;
2617
2618 case V_028714_SPI_SHADER_UINT16_ABGR:
2619 args->enabled_channels = 0x5;
2620 packi = ac_build_cvt_pk_u16;
2621 if (is_16bit) {
2622 for (unsigned chan = 0; chan < 4; chan++)
2623 values[chan] = LLVMBuildZExt(ctx->ac.builder,
2624 ac_to_integer(&ctx->ac, values[chan]),
2625 ctx->ac.i32, "");
2626 }
2627 break;
2628
2629 case V_028714_SPI_SHADER_SINT16_ABGR:
2630 args->enabled_channels = 0x5;
2631 packi = ac_build_cvt_pk_i16;
2632 if (is_16bit) {
2633 for (unsigned chan = 0; chan < 4; chan++)
2634 values[chan] = LLVMBuildSExt(ctx->ac.builder,
2635 ac_to_integer(&ctx->ac, values[chan]),
2636 ctx->ac.i32, "");
2637 }
2638 break;
2639
2640 default:
2641 case V_028714_SPI_SHADER_32_ABGR:
2642 memcpy(&args->out[0], values, sizeof(values[0]) * 4);
2643 break;
2644 }
2645
2646 /* Pack f16 or norm_i16/u16. */
2647 if (packf) {
2648 for (chan = 0; chan < 2; chan++) {
2649 LLVMValueRef pack_args[2] = {
2650 values[2 * chan],
2651 values[2 * chan + 1]
2652 };
2653 LLVMValueRef packed;
2654
2655 packed = packf(&ctx->ac, pack_args);
2656 args->out[chan] = ac_to_float(&ctx->ac, packed);
2657 }
2658 args->compr = 1; /* COMPR flag */
2659 }
2660
2661 /* Pack i16/u16. */
2662 if (packi) {
2663 for (chan = 0; chan < 2; chan++) {
2664 LLVMValueRef pack_args[2] = {
2665 ac_to_integer(&ctx->ac, values[2 * chan]),
2666 ac_to_integer(&ctx->ac, values[2 * chan + 1])
2667 };
2668 LLVMValueRef packed;
2669
2670 packed = packi(&ctx->ac, pack_args,
2671 is_int8 ? 8 : is_int10 ? 10 : 16,
2672 chan == 1);
2673 args->out[chan] = ac_to_float(&ctx->ac, packed);
2674 }
2675 args->compr = 1; /* COMPR flag */
2676 }
2677 return;
2678 }
2679
2680 if (is_16bit) {
2681 for (unsigned chan = 0; chan < 4; chan++) {
2682 values[chan] = LLVMBuildBitCast(ctx->ac.builder, values[chan], ctx->ac.i16, "");
2683 args->out[chan] = LLVMBuildZExt(ctx->ac.builder, values[chan], ctx->ac.i32, "");
2684 }
2685 } else
2686 memcpy(&args->out[0], values, sizeof(values[0]) * 4);
2687
2688 for (unsigned i = 0; i < 4; ++i)
2689 args->out[i] = ac_to_float(&ctx->ac, args->out[i]);
2690 }
2691
2692 static void
2693 radv_export_param(struct radv_shader_context *ctx, unsigned index,
2694 LLVMValueRef *values, unsigned enabled_channels)
2695 {
2696 struct ac_export_args args;
2697
2698 si_llvm_init_export_args(ctx, values, enabled_channels,
2699 V_008DFC_SQ_EXP_PARAM + index, &args);
2700 ac_build_export(&ctx->ac, &args);
2701 }
2702
2703 static LLVMValueRef
2704 radv_load_output(struct radv_shader_context *ctx, unsigned index, unsigned chan)
2705 {
2706 LLVMValueRef output =
2707 ctx->abi.outputs[ac_llvm_reg_index_soa(index, chan)];
2708
2709 return LLVMBuildLoad(ctx->ac.builder, output, "");
2710 }
2711
2712 static void
2713 radv_emit_stream_output(struct radv_shader_context *ctx,
2714 LLVMValueRef const *so_buffers,
2715 LLVMValueRef const *so_write_offsets,
2716 const struct radv_stream_output *output)
2717 {
2718 unsigned num_comps = util_bitcount(output->component_mask);
2719 unsigned loc = output->location;
2720 unsigned buf = output->buffer;
2721 unsigned offset = output->offset;
2722 unsigned start;
2723 LLVMValueRef out[4];
2724
2725 assert(num_comps && num_comps <= 4);
2726 if (!num_comps || num_comps > 4)
2727 return;
2728
2729 /* Get the first component. */
2730 start = ffs(output->component_mask) - 1;
2731
2732 /* Load the output as int. */
2733 for (int i = 0; i < num_comps; i++) {
2734 out[i] = ac_to_integer(&ctx->ac,
2735 radv_load_output(ctx, loc, start + i));
2736 }
2737
2738 /* Pack the output. */
2739 LLVMValueRef vdata = NULL;
2740
2741 switch (num_comps) {
2742 case 1: /* as i32 */
2743 vdata = out[0];
2744 break;
2745 case 2: /* as v2i32 */
2746 case 3: /* as v4i32 (aligned to 4) */
2747 out[3] = LLVMGetUndef(ctx->ac.i32);
2748 /* fall through */
2749 case 4: /* as v4i32 */
2750 vdata = ac_build_gather_values(&ctx->ac, out,
2751 util_next_power_of_two(num_comps));
2752 break;
2753 }
2754
2755 ac_build_buffer_store_dword(&ctx->ac, so_buffers[buf],
2756 vdata, num_comps, so_write_offsets[buf],
2757 ctx->ac.i32_0, offset,
2758 1, 1, true, false);
2759 }
2760
2761 static void
2762 radv_emit_streamout(struct radv_shader_context *ctx, unsigned stream)
2763 {
2764 struct ac_build_if_state if_ctx;
2765 int i;
2766
2767 /* Get bits [22:16], i.e. (so_param >> 16) & 127; */
2768 assert(ctx->streamout_config);
2769 LLVMValueRef so_vtx_count =
2770 ac_build_bfe(&ctx->ac, ctx->streamout_config,
2771 LLVMConstInt(ctx->ac.i32, 16, false),
2772 LLVMConstInt(ctx->ac.i32, 7, false), false);
2773
2774 LLVMValueRef tid = ac_get_thread_id(&ctx->ac);
2775
2776 /* can_emit = tid < so_vtx_count; */
2777 LLVMValueRef can_emit = LLVMBuildICmp(ctx->ac.builder, LLVMIntULT,
2778 tid, so_vtx_count, "");
2779
2780 /* Emit the streamout code conditionally. This actually avoids
2781 * out-of-bounds buffer access. The hw tells us via the SGPR
2782 * (so_vtx_count) which threads are allowed to emit streamout data.
2783 */
2784 ac_nir_build_if(&if_ctx, ctx, can_emit);
2785 {
2786 /* The buffer offset is computed as follows:
2787 * ByteOffset = streamout_offset[buffer_id]*4 +
2788 * (streamout_write_index + thread_id)*stride[buffer_id] +
2789 * attrib_offset
2790 */
2791 LLVMValueRef so_write_index = ctx->streamout_write_idx;
2792
2793 /* Compute (streamout_write_index + thread_id). */
2794 so_write_index =
2795 LLVMBuildAdd(ctx->ac.builder, so_write_index, tid, "");
2796
2797 /* Load the descriptor and compute the write offset for each
2798 * enabled buffer.
2799 */
2800 LLVMValueRef so_write_offset[4] = {};
2801 LLVMValueRef so_buffers[4] = {};
2802 LLVMValueRef buf_ptr = ctx->streamout_buffers;
2803
2804 for (i = 0; i < 4; i++) {
2805 uint16_t stride = ctx->shader_info->info.so.strides[i];
2806
2807 if (!stride)
2808 continue;
2809
2810 LLVMValueRef offset =
2811 LLVMConstInt(ctx->ac.i32, i, false);
2812
2813 so_buffers[i] = ac_build_load_to_sgpr(&ctx->ac,
2814 buf_ptr, offset);
2815
2816 LLVMValueRef so_offset = ctx->streamout_offset[i];
2817
2818 so_offset = LLVMBuildMul(ctx->ac.builder, so_offset,
2819 LLVMConstInt(ctx->ac.i32, 4, false), "");
2820
2821 so_write_offset[i] =
2822 ac_build_imad(&ctx->ac, so_write_index,
2823 LLVMConstInt(ctx->ac.i32,
2824 stride * 4, false),
2825 so_offset);
2826 }
2827
2828 /* Write streamout data. */
2829 for (i = 0; i < ctx->shader_info->info.so.num_outputs; i++) {
2830 struct radv_stream_output *output =
2831 &ctx->shader_info->info.so.outputs[i];
2832
2833 if (stream != output->stream)
2834 continue;
2835
2836 radv_emit_stream_output(ctx, so_buffers,
2837 so_write_offset, output);
2838 }
2839 }
2840 ac_nir_build_endif(&if_ctx);
2841 }
2842
2843 static void
2844 handle_vs_outputs_post(struct radv_shader_context *ctx,
2845 bool export_prim_id, bool export_layer_id,
2846 struct radv_vs_output_info *outinfo)
2847 {
2848 uint32_t param_count = 0;
2849 unsigned target;
2850 unsigned pos_idx, num_pos_exports = 0;
2851 struct ac_export_args args, pos_args[4] = {};
2852 LLVMValueRef psize_value = NULL, layer_value = NULL, viewport_index_value = NULL;
2853 int i;
2854
2855 if (ctx->options->key.has_multiview_view_index) {
2856 LLVMValueRef* tmp_out = &ctx->abi.outputs[ac_llvm_reg_index_soa(VARYING_SLOT_LAYER, 0)];
2857 if(!*tmp_out) {
2858 for(unsigned i = 0; i < 4; ++i)
2859 ctx->abi.outputs[ac_llvm_reg_index_soa(VARYING_SLOT_LAYER, i)] =
2860 ac_build_alloca_undef(&ctx->ac, ctx->ac.f32, "");
2861 }
2862
2863 LLVMBuildStore(ctx->ac.builder, ac_to_float(&ctx->ac, ctx->abi.view_index), *tmp_out);
2864 ctx->output_mask |= 1ull << VARYING_SLOT_LAYER;
2865 }
2866
2867 memset(outinfo->vs_output_param_offset, AC_EXP_PARAM_UNDEFINED,
2868 sizeof(outinfo->vs_output_param_offset));
2869
2870 for(unsigned location = VARYING_SLOT_CLIP_DIST0; location <= VARYING_SLOT_CLIP_DIST1; ++location) {
2871 if (ctx->output_mask & (1ull << location)) {
2872 unsigned output_usage_mask, length;
2873 LLVMValueRef slots[4];
2874 unsigned j;
2875
2876 if (ctx->stage == MESA_SHADER_VERTEX &&
2877 !ctx->is_gs_copy_shader) {
2878 output_usage_mask =
2879 ctx->shader_info->info.vs.output_usage_mask[location];
2880 } else if (ctx->stage == MESA_SHADER_TESS_EVAL) {
2881 output_usage_mask =
2882 ctx->shader_info->info.tes.output_usage_mask[location];
2883 } else {
2884 assert(ctx->is_gs_copy_shader);
2885 output_usage_mask =
2886 ctx->shader_info->info.gs.output_usage_mask[location];
2887 }
2888
2889 length = util_last_bit(output_usage_mask);
2890
2891 for (j = 0; j < length; j++)
2892 slots[j] = ac_to_float(&ctx->ac, radv_load_output(ctx, location, j));
2893
2894 for (i = length; i < 4; i++)
2895 slots[i] = LLVMGetUndef(ctx->ac.f32);
2896
2897 target = V_008DFC_SQ_EXP_POS + 2 + (location - VARYING_SLOT_CLIP_DIST0);
2898 si_llvm_init_export_args(ctx, &slots[0], 0xf, target, &args);
2899 memcpy(&pos_args[target - V_008DFC_SQ_EXP_POS],
2900 &args, sizeof(args));
2901
2902 /* Export the clip/cull distances values to the next stage. */
2903 radv_export_param(ctx, param_count, &slots[0], 0xf);
2904 outinfo->vs_output_param_offset[location] = param_count++;
2905 }
2906 }
2907
2908 LLVMValueRef pos_values[4] = {ctx->ac.f32_0, ctx->ac.f32_0, ctx->ac.f32_0, ctx->ac.f32_1};
2909 if (ctx->output_mask & (1ull << VARYING_SLOT_POS)) {
2910 for (unsigned j = 0; j < 4; j++)
2911 pos_values[j] = radv_load_output(ctx, VARYING_SLOT_POS, j);
2912 }
2913 si_llvm_init_export_args(ctx, pos_values, 0xf, V_008DFC_SQ_EXP_POS, &pos_args[0]);
2914
2915 if (ctx->output_mask & (1ull << VARYING_SLOT_PSIZ)) {
2916 outinfo->writes_pointsize = true;
2917 psize_value = radv_load_output(ctx, VARYING_SLOT_PSIZ, 0);
2918 }
2919
2920 if (ctx->output_mask & (1ull << VARYING_SLOT_LAYER)) {
2921 outinfo->writes_layer = true;
2922 layer_value = radv_load_output(ctx, VARYING_SLOT_LAYER, 0);
2923 }
2924
2925 if (ctx->output_mask & (1ull << VARYING_SLOT_VIEWPORT)) {
2926 outinfo->writes_viewport_index = true;
2927 viewport_index_value = radv_load_output(ctx, VARYING_SLOT_VIEWPORT, 0);
2928 }
2929
2930 if (ctx->shader_info->info.so.num_outputs &&
2931 !ctx->is_gs_copy_shader) {
2932 /* The GS copy shader emission already emits streamout. */
2933 radv_emit_streamout(ctx, 0);
2934 }
2935
2936 if (outinfo->writes_pointsize ||
2937 outinfo->writes_layer ||
2938 outinfo->writes_viewport_index) {
2939 pos_args[1].enabled_channels = ((outinfo->writes_pointsize == true ? 1 : 0) |
2940 (outinfo->writes_layer == true ? 4 : 0));
2941 pos_args[1].valid_mask = 0;
2942 pos_args[1].done = 0;
2943 pos_args[1].target = V_008DFC_SQ_EXP_POS + 1;
2944 pos_args[1].compr = 0;
2945 pos_args[1].out[0] = ctx->ac.f32_0; /* X */
2946 pos_args[1].out[1] = ctx->ac.f32_0; /* Y */
2947 pos_args[1].out[2] = ctx->ac.f32_0; /* Z */
2948 pos_args[1].out[3] = ctx->ac.f32_0; /* W */
2949
2950 if (outinfo->writes_pointsize == true)
2951 pos_args[1].out[0] = psize_value;
2952 if (outinfo->writes_layer == true)
2953 pos_args[1].out[2] = layer_value;
2954 if (outinfo->writes_viewport_index == true) {
2955 if (ctx->options->chip_class >= GFX9) {
2956 /* GFX9 has the layer in out.z[10:0] and the viewport
2957 * index in out.z[19:16].
2958 */
2959 LLVMValueRef v = viewport_index_value;
2960 v = ac_to_integer(&ctx->ac, v);
2961 v = LLVMBuildShl(ctx->ac.builder, v,
2962 LLVMConstInt(ctx->ac.i32, 16, false),
2963 "");
2964 v = LLVMBuildOr(ctx->ac.builder, v,
2965 ac_to_integer(&ctx->ac, pos_args[1].out[2]), "");
2966
2967 pos_args[1].out[2] = ac_to_float(&ctx->ac, v);
2968 pos_args[1].enabled_channels |= 1 << 2;
2969 } else {
2970 pos_args[1].out[3] = viewport_index_value;
2971 pos_args[1].enabled_channels |= 1 << 3;
2972 }
2973 }
2974 }
2975 for (i = 0; i < 4; i++) {
2976 if (pos_args[i].out[0])
2977 num_pos_exports++;
2978 }
2979
2980 pos_idx = 0;
2981 for (i = 0; i < 4; i++) {
2982 if (!pos_args[i].out[0])
2983 continue;
2984
2985 /* Specify the target we are exporting */
2986 pos_args[i].target = V_008DFC_SQ_EXP_POS + pos_idx++;
2987 if (pos_idx == num_pos_exports)
2988 pos_args[i].done = 1;
2989 ac_build_export(&ctx->ac, &pos_args[i]);
2990 }
2991
2992 for (unsigned i = 0; i < AC_LLVM_MAX_OUTPUTS; ++i) {
2993 LLVMValueRef values[4];
2994 if (!(ctx->output_mask & (1ull << i)))
2995 continue;
2996
2997 if (i != VARYING_SLOT_LAYER &&
2998 i != VARYING_SLOT_PRIMITIVE_ID &&
2999 i < VARYING_SLOT_VAR0)
3000 continue;
3001
3002 for (unsigned j = 0; j < 4; j++)
3003 values[j] = ac_to_float(&ctx->ac, radv_load_output(ctx, i, j));
3004
3005 unsigned output_usage_mask;
3006
3007 if (ctx->stage == MESA_SHADER_VERTEX &&
3008 !ctx->is_gs_copy_shader) {
3009 output_usage_mask =
3010 ctx->shader_info->info.vs.output_usage_mask[i];
3011 } else if (ctx->stage == MESA_SHADER_TESS_EVAL) {
3012 output_usage_mask =
3013 ctx->shader_info->info.tes.output_usage_mask[i];
3014 } else {
3015 assert(ctx->is_gs_copy_shader);
3016 output_usage_mask =
3017 ctx->shader_info->info.gs.output_usage_mask[i];
3018 }
3019
3020 radv_export_param(ctx, param_count, values, output_usage_mask);
3021
3022 outinfo->vs_output_param_offset[i] = param_count++;
3023 }
3024
3025 if (export_prim_id) {
3026 LLVMValueRef values[4];
3027
3028 values[0] = ctx->vs_prim_id;
3029 ctx->shader_info->vs.vgpr_comp_cnt = MAX2(2,
3030 ctx->shader_info->vs.vgpr_comp_cnt);
3031 for (unsigned j = 1; j < 4; j++)
3032 values[j] = ctx->ac.f32_0;
3033
3034 radv_export_param(ctx, param_count, values, 0x1);
3035
3036 outinfo->vs_output_param_offset[VARYING_SLOT_PRIMITIVE_ID] = param_count++;
3037 outinfo->export_prim_id = true;
3038 }
3039
3040 if (export_layer_id && layer_value) {
3041 LLVMValueRef values[4];
3042
3043 values[0] = layer_value;
3044 for (unsigned j = 1; j < 4; j++)
3045 values[j] = ctx->ac.f32_0;
3046
3047 radv_export_param(ctx, param_count, values, 0x1);
3048
3049 outinfo->vs_output_param_offset[VARYING_SLOT_LAYER] = param_count++;
3050 }
3051
3052 outinfo->pos_exports = num_pos_exports;
3053 outinfo->param_exports = param_count;
3054 }
3055
3056 static void
3057 handle_es_outputs_post(struct radv_shader_context *ctx,
3058 struct radv_es_output_info *outinfo)
3059 {
3060 int j;
3061 uint64_t max_output_written = 0;
3062 LLVMValueRef lds_base = NULL;
3063
3064 for (unsigned i = 0; i < AC_LLVM_MAX_OUTPUTS; ++i) {
3065 int param_index;
3066
3067 if (!(ctx->output_mask & (1ull << i)))
3068 continue;
3069
3070 param_index = shader_io_get_unique_index(i);
3071
3072 max_output_written = MAX2(param_index, max_output_written);
3073 }
3074
3075 outinfo->esgs_itemsize = (max_output_written + 1) * 16;
3076
3077 if (ctx->ac.chip_class >= GFX9) {
3078 unsigned itemsize_dw = outinfo->esgs_itemsize / 4;
3079 LLVMValueRef vertex_idx = ac_get_thread_id(&ctx->ac);
3080 LLVMValueRef wave_idx = ac_unpack_param(&ctx->ac, ctx->merged_wave_info, 24, 4);
3081 vertex_idx = LLVMBuildOr(ctx->ac.builder, vertex_idx,
3082 LLVMBuildMul(ctx->ac.builder, wave_idx,
3083 LLVMConstInt(ctx->ac.i32, 64, false), ""), "");
3084 lds_base = LLVMBuildMul(ctx->ac.builder, vertex_idx,
3085 LLVMConstInt(ctx->ac.i32, itemsize_dw, 0), "");
3086 }
3087
3088 for (unsigned i = 0; i < AC_LLVM_MAX_OUTPUTS; ++i) {
3089 LLVMValueRef dw_addr = NULL;
3090 LLVMValueRef *out_ptr = &ctx->abi.outputs[i * 4];
3091 unsigned output_usage_mask;
3092 int param_index;
3093
3094 if (!(ctx->output_mask & (1ull << i)))
3095 continue;
3096
3097 if (ctx->stage == MESA_SHADER_VERTEX) {
3098 output_usage_mask =
3099 ctx->shader_info->info.vs.output_usage_mask[i];
3100 } else {
3101 assert(ctx->stage == MESA_SHADER_TESS_EVAL);
3102 output_usage_mask =
3103 ctx->shader_info->info.tes.output_usage_mask[i];
3104 }
3105
3106 param_index = shader_io_get_unique_index(i);
3107
3108 if (lds_base) {
3109 dw_addr = LLVMBuildAdd(ctx->ac.builder, lds_base,
3110 LLVMConstInt(ctx->ac.i32, param_index * 4, false),
3111 "");
3112 }
3113
3114 for (j = 0; j < 4; j++) {
3115 if (!(output_usage_mask & (1 << j)))
3116 continue;
3117
3118 LLVMValueRef out_val = LLVMBuildLoad(ctx->ac.builder, out_ptr[j], "");
3119 out_val = ac_to_integer(&ctx->ac, out_val);
3120 out_val = LLVMBuildZExtOrBitCast(ctx->ac.builder, out_val, ctx->ac.i32, "");
3121
3122 if (ctx->ac.chip_class >= GFX9) {
3123 LLVMValueRef dw_addr_offset =
3124 LLVMBuildAdd(ctx->ac.builder, dw_addr,
3125 LLVMConstInt(ctx->ac.i32,
3126 j, false), "");
3127
3128 ac_lds_store(&ctx->ac, dw_addr_offset, out_val);
3129 } else {
3130 ac_build_buffer_store_dword(&ctx->ac,
3131 ctx->esgs_ring,
3132 out_val, 1,
3133 NULL, ctx->es2gs_offset,
3134 (4 * param_index + j) * 4,
3135 1, 1, true, true);
3136 }
3137 }
3138 }
3139 }
3140
3141 static void
3142 handle_ls_outputs_post(struct radv_shader_context *ctx)
3143 {
3144 LLVMValueRef vertex_id = ctx->rel_auto_id;
3145 uint32_t num_tcs_inputs = util_last_bit64(ctx->shader_info->info.vs.ls_outputs_written);
3146 LLVMValueRef vertex_dw_stride = LLVMConstInt(ctx->ac.i32, num_tcs_inputs * 4, false);
3147 LLVMValueRef base_dw_addr = LLVMBuildMul(ctx->ac.builder, vertex_id,
3148 vertex_dw_stride, "");
3149
3150 for (unsigned i = 0; i < AC_LLVM_MAX_OUTPUTS; ++i) {
3151 LLVMValueRef *out_ptr = &ctx->abi.outputs[i * 4];
3152
3153 if (!(ctx->output_mask & (1ull << i)))
3154 continue;
3155
3156 int param = shader_io_get_unique_index(i);
3157 LLVMValueRef dw_addr = LLVMBuildAdd(ctx->ac.builder, base_dw_addr,
3158 LLVMConstInt(ctx->ac.i32, param * 4, false),
3159 "");
3160 for (unsigned j = 0; j < 4; j++) {
3161 LLVMValueRef value = LLVMBuildLoad(ctx->ac.builder, out_ptr[j], "");
3162 value = ac_to_integer(&ctx->ac, value);
3163 value = LLVMBuildZExtOrBitCast(ctx->ac.builder, value, ctx->ac.i32, "");
3164 ac_lds_store(&ctx->ac, dw_addr, value);
3165 dw_addr = LLVMBuildAdd(ctx->ac.builder, dw_addr, ctx->ac.i32_1, "");
3166 }
3167 }
3168 }
3169
3170 static void
3171 write_tess_factors(struct radv_shader_context *ctx)
3172 {
3173 unsigned stride, outer_comps, inner_comps;
3174 struct ac_build_if_state if_ctx, inner_if_ctx;
3175 LLVMValueRef invocation_id = ac_unpack_param(&ctx->ac, ctx->abi.tcs_rel_ids, 8, 5);
3176 LLVMValueRef rel_patch_id = ac_unpack_param(&ctx->ac, ctx->abi.tcs_rel_ids, 0, 8);
3177 unsigned tess_inner_index = 0, tess_outer_index;
3178 LLVMValueRef lds_base, lds_inner = NULL, lds_outer, byteoffset, buffer;
3179 LLVMValueRef out[6], vec0, vec1, tf_base, inner[4], outer[4];
3180 int i;
3181 ac_emit_barrier(&ctx->ac, ctx->stage);
3182
3183 switch (ctx->options->key.tcs.primitive_mode) {
3184 case GL_ISOLINES:
3185 stride = 2;
3186 outer_comps = 2;
3187 inner_comps = 0;
3188 break;
3189 case GL_TRIANGLES:
3190 stride = 4;
3191 outer_comps = 3;
3192 inner_comps = 1;
3193 break;
3194 case GL_QUADS:
3195 stride = 6;
3196 outer_comps = 4;
3197 inner_comps = 2;
3198 break;
3199 default:
3200 return;
3201 }
3202
3203 ac_nir_build_if(&if_ctx, ctx,
3204 LLVMBuildICmp(ctx->ac.builder, LLVMIntEQ,
3205 invocation_id, ctx->ac.i32_0, ""));
3206
3207 lds_base = get_tcs_out_current_patch_data_offset(ctx);
3208
3209 if (inner_comps) {
3210 tess_inner_index = shader_io_get_unique_index(VARYING_SLOT_TESS_LEVEL_INNER);
3211 lds_inner = LLVMBuildAdd(ctx->ac.builder, lds_base,
3212 LLVMConstInt(ctx->ac.i32, tess_inner_index * 4, false), "");
3213 }
3214
3215 tess_outer_index = shader_io_get_unique_index(VARYING_SLOT_TESS_LEVEL_OUTER);
3216 lds_outer = LLVMBuildAdd(ctx->ac.builder, lds_base,
3217 LLVMConstInt(ctx->ac.i32, tess_outer_index * 4, false), "");
3218
3219 for (i = 0; i < 4; i++) {
3220 inner[i] = LLVMGetUndef(ctx->ac.i32);
3221 outer[i] = LLVMGetUndef(ctx->ac.i32);
3222 }
3223
3224 // LINES reversal
3225 if (ctx->options->key.tcs.primitive_mode == GL_ISOLINES) {
3226 outer[0] = out[1] = ac_lds_load(&ctx->ac, lds_outer);
3227 lds_outer = LLVMBuildAdd(ctx->ac.builder, lds_outer,
3228 ctx->ac.i32_1, "");
3229 outer[1] = out[0] = ac_lds_load(&ctx->ac, lds_outer);
3230 } else {
3231 for (i = 0; i < outer_comps; i++) {
3232 outer[i] = out[i] =
3233 ac_lds_load(&ctx->ac, lds_outer);
3234 lds_outer = LLVMBuildAdd(ctx->ac.builder, lds_outer,
3235 ctx->ac.i32_1, "");
3236 }
3237 for (i = 0; i < inner_comps; i++) {
3238 inner[i] = out[outer_comps+i] =
3239 ac_lds_load(&ctx->ac, lds_inner);
3240 lds_inner = LLVMBuildAdd(ctx->ac.builder, lds_inner,
3241 ctx->ac.i32_1, "");
3242 }
3243 }
3244
3245 /* Convert the outputs to vectors for stores. */
3246 vec0 = ac_build_gather_values(&ctx->ac, out, MIN2(stride, 4));
3247 vec1 = NULL;
3248
3249 if (stride > 4)
3250 vec1 = ac_build_gather_values(&ctx->ac, out + 4, stride - 4);
3251
3252
3253 buffer = ctx->hs_ring_tess_factor;
3254 tf_base = ctx->tess_factor_offset;
3255 byteoffset = LLVMBuildMul(ctx->ac.builder, rel_patch_id,
3256 LLVMConstInt(ctx->ac.i32, 4 * stride, false), "");
3257 unsigned tf_offset = 0;
3258
3259 if (ctx->options->chip_class <= VI) {
3260 ac_nir_build_if(&inner_if_ctx, ctx,
3261 LLVMBuildICmp(ctx->ac.builder, LLVMIntEQ,
3262 rel_patch_id, ctx->ac.i32_0, ""));
3263
3264 /* Store the dynamic HS control word. */
3265 ac_build_buffer_store_dword(&ctx->ac, buffer,
3266 LLVMConstInt(ctx->ac.i32, 0x80000000, false),
3267 1, ctx->ac.i32_0, tf_base,
3268 0, 1, 0, true, false);
3269 tf_offset += 4;
3270
3271 ac_nir_build_endif(&inner_if_ctx);
3272 }
3273
3274 /* Store the tessellation factors. */
3275 ac_build_buffer_store_dword(&ctx->ac, buffer, vec0,
3276 MIN2(stride, 4), byteoffset, tf_base,
3277 tf_offset, 1, 0, true, false);
3278 if (vec1)
3279 ac_build_buffer_store_dword(&ctx->ac, buffer, vec1,
3280 stride - 4, byteoffset, tf_base,
3281 16 + tf_offset, 1, 0, true, false);
3282
3283 //store to offchip for TES to read - only if TES reads them
3284 if (ctx->options->key.tcs.tes_reads_tess_factors) {
3285 LLVMValueRef inner_vec, outer_vec, tf_outer_offset;
3286 LLVMValueRef tf_inner_offset;
3287 unsigned param_outer, param_inner;
3288
3289 param_outer = shader_io_get_unique_index(VARYING_SLOT_TESS_LEVEL_OUTER);
3290 tf_outer_offset = get_tcs_tes_buffer_address(ctx, NULL,
3291 LLVMConstInt(ctx->ac.i32, param_outer, 0));
3292
3293 outer_vec = ac_build_gather_values(&ctx->ac, outer,
3294 util_next_power_of_two(outer_comps));
3295
3296 ac_build_buffer_store_dword(&ctx->ac, ctx->hs_ring_tess_offchip, outer_vec,
3297 outer_comps, tf_outer_offset,
3298 ctx->oc_lds, 0, 1, 0, true, false);
3299 if (inner_comps) {
3300 param_inner = shader_io_get_unique_index(VARYING_SLOT_TESS_LEVEL_INNER);
3301 tf_inner_offset = get_tcs_tes_buffer_address(ctx, NULL,
3302 LLVMConstInt(ctx->ac.i32, param_inner, 0));
3303
3304 inner_vec = inner_comps == 1 ? inner[0] :
3305 ac_build_gather_values(&ctx->ac, inner, inner_comps);
3306 ac_build_buffer_store_dword(&ctx->ac, ctx->hs_ring_tess_offchip, inner_vec,
3307 inner_comps, tf_inner_offset,
3308 ctx->oc_lds, 0, 1, 0, true, false);
3309 }
3310 }
3311 ac_nir_build_endif(&if_ctx);
3312 }
3313
3314 static void
3315 handle_tcs_outputs_post(struct radv_shader_context *ctx)
3316 {
3317 write_tess_factors(ctx);
3318 }
3319
3320 static bool
3321 si_export_mrt_color(struct radv_shader_context *ctx,
3322 LLVMValueRef *color, unsigned index,
3323 struct ac_export_args *args)
3324 {
3325 /* Export */
3326 si_llvm_init_export_args(ctx, color, 0xf,
3327 V_008DFC_SQ_EXP_MRT + index, args);
3328 if (!args->enabled_channels)
3329 return false; /* unnecessary NULL export */
3330
3331 return true;
3332 }
3333
3334 static void
3335 radv_export_mrt_z(struct radv_shader_context *ctx,
3336 LLVMValueRef depth, LLVMValueRef stencil,
3337 LLVMValueRef samplemask)
3338 {
3339 struct ac_export_args args;
3340
3341 ac_export_mrt_z(&ctx->ac, depth, stencil, samplemask, &args);
3342
3343 ac_build_export(&ctx->ac, &args);
3344 }
3345
3346 static void
3347 handle_fs_outputs_post(struct radv_shader_context *ctx)
3348 {
3349 unsigned index = 0;
3350 LLVMValueRef depth = NULL, stencil = NULL, samplemask = NULL;
3351 struct ac_export_args color_args[8];
3352
3353 for (unsigned i = 0; i < AC_LLVM_MAX_OUTPUTS; ++i) {
3354 LLVMValueRef values[4];
3355
3356 if (!(ctx->output_mask & (1ull << i)))
3357 continue;
3358
3359 if (i < FRAG_RESULT_DATA0)
3360 continue;
3361
3362 for (unsigned j = 0; j < 4; j++)
3363 values[j] = ac_to_float(&ctx->ac,
3364 radv_load_output(ctx, i, j));
3365
3366 bool ret = si_export_mrt_color(ctx, values,
3367 i - FRAG_RESULT_DATA0,
3368 &color_args[index]);
3369 if (ret)
3370 index++;
3371 }
3372
3373 /* Process depth, stencil, samplemask. */
3374 if (ctx->shader_info->info.ps.writes_z) {
3375 depth = ac_to_float(&ctx->ac,
3376 radv_load_output(ctx, FRAG_RESULT_DEPTH, 0));
3377 }
3378 if (ctx->shader_info->info.ps.writes_stencil) {
3379 stencil = ac_to_float(&ctx->ac,
3380 radv_load_output(ctx, FRAG_RESULT_STENCIL, 0));
3381 }
3382 if (ctx->shader_info->info.ps.writes_sample_mask) {
3383 samplemask = ac_to_float(&ctx->ac,
3384 radv_load_output(ctx, FRAG_RESULT_SAMPLE_MASK, 0));
3385 }
3386
3387 /* Set the DONE bit on last non-null color export only if Z isn't
3388 * exported.
3389 */
3390 if (index > 0 &&
3391 !ctx->shader_info->info.ps.writes_z &&
3392 !ctx->shader_info->info.ps.writes_stencil &&
3393 !ctx->shader_info->info.ps.writes_sample_mask) {
3394 unsigned last = index - 1;
3395
3396 color_args[last].valid_mask = 1; /* whether the EXEC mask is valid */
3397 color_args[last].done = 1; /* DONE bit */
3398 }
3399
3400 /* Export PS outputs. */
3401 for (unsigned i = 0; i < index; i++)
3402 ac_build_export(&ctx->ac, &color_args[i]);
3403
3404 if (depth || stencil || samplemask)
3405 radv_export_mrt_z(ctx, depth, stencil, samplemask);
3406 else if (!index)
3407 ac_build_export_null(&ctx->ac);
3408 }
3409
3410 static void
3411 emit_gs_epilogue(struct radv_shader_context *ctx)
3412 {
3413 ac_build_sendmsg(&ctx->ac, AC_SENDMSG_GS_OP_NOP | AC_SENDMSG_GS_DONE, ctx->gs_wave_id);
3414 }
3415
3416 static void
3417 handle_shader_outputs_post(struct ac_shader_abi *abi, unsigned max_outputs,
3418 LLVMValueRef *addrs)
3419 {
3420 struct radv_shader_context *ctx = radv_shader_context_from_abi(abi);
3421
3422 switch (ctx->stage) {
3423 case MESA_SHADER_VERTEX:
3424 if (ctx->options->key.vs.as_ls)
3425 handle_ls_outputs_post(ctx);
3426 else if (ctx->options->key.vs.as_es)
3427 handle_es_outputs_post(ctx, &ctx->shader_info->vs.es_info);
3428 else
3429 handle_vs_outputs_post(ctx, ctx->options->key.vs.export_prim_id,
3430 ctx->options->key.vs.export_layer_id,
3431 &ctx->shader_info->vs.outinfo);
3432 break;
3433 case MESA_SHADER_FRAGMENT:
3434 handle_fs_outputs_post(ctx);
3435 break;
3436 case MESA_SHADER_GEOMETRY:
3437 emit_gs_epilogue(ctx);
3438 break;
3439 case MESA_SHADER_TESS_CTRL:
3440 handle_tcs_outputs_post(ctx);
3441 break;
3442 case MESA_SHADER_TESS_EVAL:
3443 if (ctx->options->key.tes.as_es)
3444 handle_es_outputs_post(ctx, &ctx->shader_info->tes.es_info);
3445 else
3446 handle_vs_outputs_post(ctx, ctx->options->key.tes.export_prim_id,
3447 ctx->options->key.tes.export_layer_id,
3448 &ctx->shader_info->tes.outinfo);
3449 break;
3450 default:
3451 break;
3452 }
3453 }
3454
3455 static void ac_llvm_finalize_module(struct radv_shader_context *ctx,
3456 LLVMPassManagerRef passmgr,
3457 const struct radv_nir_compiler_options *options)
3458 {
3459 LLVMRunPassManager(passmgr, ctx->ac.module);
3460 LLVMDisposeBuilder(ctx->ac.builder);
3461
3462 ac_llvm_context_dispose(&ctx->ac);
3463 }
3464
3465 static void
3466 ac_nir_eliminate_const_vs_outputs(struct radv_shader_context *ctx)
3467 {
3468 struct radv_vs_output_info *outinfo;
3469
3470 switch (ctx->stage) {
3471 case MESA_SHADER_FRAGMENT:
3472 case MESA_SHADER_COMPUTE:
3473 case MESA_SHADER_TESS_CTRL:
3474 case MESA_SHADER_GEOMETRY:
3475 return;
3476 case MESA_SHADER_VERTEX:
3477 if (ctx->options->key.vs.as_ls ||
3478 ctx->options->key.vs.as_es)
3479 return;
3480 outinfo = &ctx->shader_info->vs.outinfo;
3481 break;
3482 case MESA_SHADER_TESS_EVAL:
3483 if (ctx->options->key.vs.as_es)
3484 return;
3485 outinfo = &ctx->shader_info->tes.outinfo;
3486 break;
3487 default:
3488 unreachable("Unhandled shader type");
3489 }
3490
3491 ac_optimize_vs_outputs(&ctx->ac,
3492 ctx->main_function,
3493 outinfo->vs_output_param_offset,
3494 VARYING_SLOT_MAX,
3495 &outinfo->param_exports);
3496 }
3497
3498 static void
3499 ac_setup_rings(struct radv_shader_context *ctx)
3500 {
3501 if (ctx->options->chip_class <= VI &&
3502 (ctx->stage == MESA_SHADER_GEOMETRY ||
3503 ctx->options->key.vs.as_es || ctx->options->key.tes.as_es)) {
3504 unsigned ring = ctx->stage == MESA_SHADER_GEOMETRY ? RING_ESGS_GS
3505 : RING_ESGS_VS;
3506 LLVMValueRef offset = LLVMConstInt(ctx->ac.i32, ring, false);
3507
3508 ctx->esgs_ring = ac_build_load_to_sgpr(&ctx->ac,
3509 ctx->ring_offsets,
3510 offset);
3511 }
3512
3513 if (ctx->is_gs_copy_shader) {
3514 ctx->gsvs_ring[0] =
3515 ac_build_load_to_sgpr(&ctx->ac, ctx->ring_offsets,
3516 LLVMConstInt(ctx->ac.i32,
3517 RING_GSVS_VS, false));
3518 }
3519
3520 if (ctx->stage == MESA_SHADER_GEOMETRY) {
3521 /* The conceptual layout of the GSVS ring is
3522 * v0c0 .. vLv0 v0c1 .. vLc1 ..
3523 * but the real memory layout is swizzled across
3524 * threads:
3525 * t0v0c0 .. t15v0c0 t0v1c0 .. t15v1c0 ... t15vLcL
3526 * t16v0c0 ..
3527 * Override the buffer descriptor accordingly.
3528 */
3529 LLVMTypeRef v2i64 = LLVMVectorType(ctx->ac.i64, 2);
3530 uint64_t stream_offset = 0;
3531 unsigned num_records = 64;
3532 LLVMValueRef base_ring;
3533
3534 base_ring =
3535 ac_build_load_to_sgpr(&ctx->ac, ctx->ring_offsets,
3536 LLVMConstInt(ctx->ac.i32,
3537 RING_GSVS_GS, false));
3538
3539 for (unsigned stream = 0; stream < 4; stream++) {
3540 unsigned num_components, stride;
3541 LLVMValueRef ring, tmp;
3542
3543 num_components =
3544 ctx->shader_info->info.gs.num_stream_output_components[stream];
3545
3546 if (!num_components)
3547 continue;
3548
3549 stride = 4 * num_components * ctx->gs_max_out_vertices;
3550
3551 /* Limit on the stride field for <= CIK. */
3552 assert(stride < (1 << 14));
3553
3554 ring = LLVMBuildBitCast(ctx->ac.builder,
3555 base_ring, v2i64, "");
3556 tmp = LLVMBuildExtractElement(ctx->ac.builder,
3557 ring, ctx->ac.i32_0, "");
3558 tmp = LLVMBuildAdd(ctx->ac.builder, tmp,
3559 LLVMConstInt(ctx->ac.i64,
3560 stream_offset, 0), "");
3561 ring = LLVMBuildInsertElement(ctx->ac.builder,
3562 ring, tmp, ctx->ac.i32_0, "");
3563
3564 stream_offset += stride * 64;
3565
3566 ring = LLVMBuildBitCast(ctx->ac.builder, ring,
3567 ctx->ac.v4i32, "");
3568
3569 tmp = LLVMBuildExtractElement(ctx->ac.builder, ring,
3570 ctx->ac.i32_1, "");
3571 tmp = LLVMBuildOr(ctx->ac.builder, tmp,
3572 LLVMConstInt(ctx->ac.i32,
3573 S_008F04_STRIDE(stride), false), "");
3574 ring = LLVMBuildInsertElement(ctx->ac.builder, ring, tmp,
3575 ctx->ac.i32_1, "");
3576
3577 ring = LLVMBuildInsertElement(ctx->ac.builder, ring,
3578 LLVMConstInt(ctx->ac.i32,
3579 num_records, false),
3580 LLVMConstInt(ctx->ac.i32, 2, false), "");
3581
3582 ctx->gsvs_ring[stream] = ring;
3583 }
3584 }
3585
3586 if (ctx->stage == MESA_SHADER_TESS_CTRL ||
3587 ctx->stage == MESA_SHADER_TESS_EVAL) {
3588 ctx->hs_ring_tess_offchip = ac_build_load_to_sgpr(&ctx->ac, ctx->ring_offsets, LLVMConstInt(ctx->ac.i32, RING_HS_TESS_OFFCHIP, false));
3589 ctx->hs_ring_tess_factor = ac_build_load_to_sgpr(&ctx->ac, ctx->ring_offsets, LLVMConstInt(ctx->ac.i32, RING_HS_TESS_FACTOR, false));
3590 }
3591 }
3592
3593 unsigned
3594 radv_nir_get_max_workgroup_size(enum chip_class chip_class,
3595 const struct nir_shader *nir)
3596 {
3597 switch (nir->info.stage) {
3598 case MESA_SHADER_TESS_CTRL:
3599 return chip_class >= CIK ? 128 : 64;
3600 case MESA_SHADER_GEOMETRY:
3601 return chip_class >= GFX9 ? 128 : 64;
3602 case MESA_SHADER_COMPUTE:
3603 break;
3604 default:
3605 return 0;
3606 }
3607
3608 unsigned max_workgroup_size = nir->info.cs.local_size[0] *
3609 nir->info.cs.local_size[1] *
3610 nir->info.cs.local_size[2];
3611 return max_workgroup_size;
3612 }
3613
3614 /* Fixup the HW not emitting the TCS regs if there are no HS threads. */
3615 static void ac_nir_fixup_ls_hs_input_vgprs(struct radv_shader_context *ctx)
3616 {
3617 LLVMValueRef count = ac_unpack_param(&ctx->ac, ctx->merged_wave_info, 8, 8);
3618 LLVMValueRef hs_empty = LLVMBuildICmp(ctx->ac.builder, LLVMIntEQ, count,
3619 ctx->ac.i32_0, "");
3620 ctx->abi.instance_id = LLVMBuildSelect(ctx->ac.builder, hs_empty, ctx->rel_auto_id, ctx->abi.instance_id, "");
3621 ctx->rel_auto_id = LLVMBuildSelect(ctx->ac.builder, hs_empty, ctx->abi.tcs_rel_ids, ctx->rel_auto_id, "");
3622 ctx->abi.vertex_id = LLVMBuildSelect(ctx->ac.builder, hs_empty, ctx->abi.tcs_patch_id, ctx->abi.vertex_id, "");
3623 }
3624
3625 static void prepare_gs_input_vgprs(struct radv_shader_context *ctx)
3626 {
3627 for(int i = 5; i >= 0; --i) {
3628 ctx->gs_vtx_offset[i] = ac_unpack_param(&ctx->ac, ctx->gs_vtx_offset[i & ~1],
3629 (i & 1) * 16, 16);
3630 }
3631
3632 ctx->gs_wave_id = ac_unpack_param(&ctx->ac, ctx->merged_wave_info, 16, 8);
3633 }
3634
3635
3636 static
3637 LLVMModuleRef ac_translate_nir_to_llvm(struct ac_llvm_compiler *ac_llvm,
3638 struct nir_shader *const *shaders,
3639 int shader_count,
3640 struct radv_shader_variant_info *shader_info,
3641 const struct radv_nir_compiler_options *options)
3642 {
3643 struct radv_shader_context ctx = {0};
3644 unsigned i;
3645 ctx.options = options;
3646 ctx.shader_info = shader_info;
3647
3648 ac_llvm_context_init(&ctx.ac, options->chip_class, options->family);
3649 ctx.context = ctx.ac.context;
3650 ctx.ac.module = ac_create_module(ac_llvm->tm, ctx.context);
3651
3652 enum ac_float_mode float_mode =
3653 options->unsafe_math ? AC_FLOAT_MODE_UNSAFE_FP_MATH :
3654 AC_FLOAT_MODE_DEFAULT;
3655
3656 ctx.ac.builder = ac_create_builder(ctx.context, float_mode);
3657
3658 memset(shader_info, 0, sizeof(*shader_info));
3659
3660 radv_nir_shader_info_init(&shader_info->info);
3661
3662 for(int i = 0; i < shader_count; ++i)
3663 radv_nir_shader_info_pass(shaders[i], options, &shader_info->info);
3664
3665 for (i = 0; i < RADV_UD_MAX_SETS; i++)
3666 shader_info->user_sgprs_locs.descriptor_sets[i].sgpr_idx = -1;
3667 for (i = 0; i < AC_UD_MAX_UD; i++)
3668 shader_info->user_sgprs_locs.shader_data[i].sgpr_idx = -1;
3669
3670 ctx.max_workgroup_size = 0;
3671 for (int i = 0; i < shader_count; ++i) {
3672 ctx.max_workgroup_size = MAX2(ctx.max_workgroup_size,
3673 radv_nir_get_max_workgroup_size(ctx.options->chip_class,
3674 shaders[i]));
3675 }
3676
3677 create_function(&ctx, shaders[shader_count - 1]->info.stage, shader_count >= 2,
3678 shader_count >= 2 ? shaders[shader_count - 2]->info.stage : MESA_SHADER_VERTEX);
3679
3680 ctx.abi.inputs = &ctx.inputs[0];
3681 ctx.abi.emit_outputs = handle_shader_outputs_post;
3682 ctx.abi.emit_vertex = visit_emit_vertex;
3683 ctx.abi.load_ubo = radv_load_ubo;
3684 ctx.abi.load_ssbo = radv_load_ssbo;
3685 ctx.abi.load_sampler_desc = radv_get_sampler_desc;
3686 ctx.abi.load_resource = radv_load_resource;
3687 ctx.abi.clamp_shadow_reference = false;
3688 ctx.abi.gfx9_stride_size_workaround = ctx.ac.chip_class == GFX9 && HAVE_LLVM < 0x800;
3689
3690 /* Because the new raw/struct atomic intrinsics are buggy with LLVM 8,
3691 * we fallback to the old intrinsics for atomic buffer image operations
3692 * and thus we need to apply the indexing workaround...
3693 */
3694 ctx.abi.gfx9_stride_size_workaround_for_atomic = ctx.ac.chip_class == GFX9 && HAVE_LLVM < 0x900;
3695
3696 if (shader_count >= 2)
3697 ac_init_exec_full_mask(&ctx.ac);
3698
3699 if ((ctx.ac.family == CHIP_VEGA10 ||
3700 ctx.ac.family == CHIP_RAVEN) &&
3701 shaders[shader_count - 1]->info.stage == MESA_SHADER_TESS_CTRL)
3702 ac_nir_fixup_ls_hs_input_vgprs(&ctx);
3703
3704 for(int i = 0; i < shader_count; ++i) {
3705 ctx.stage = shaders[i]->info.stage;
3706 ctx.output_mask = 0;
3707
3708 if (shaders[i]->info.stage == MESA_SHADER_GEOMETRY) {
3709 for (int i = 0; i < 4; i++) {
3710 ctx.gs_next_vertex[i] =
3711 ac_build_alloca(&ctx.ac, ctx.ac.i32, "");
3712 }
3713 ctx.gs_max_out_vertices = shaders[i]->info.gs.vertices_out;
3714 ctx.abi.load_inputs = load_gs_input;
3715 ctx.abi.emit_primitive = visit_end_primitive;
3716 } else if (shaders[i]->info.stage == MESA_SHADER_TESS_CTRL) {
3717 ctx.tcs_outputs_read = shaders[i]->info.outputs_read;
3718 ctx.tcs_patch_outputs_read = shaders[i]->info.patch_outputs_read;
3719 ctx.abi.load_tess_varyings = load_tcs_varyings;
3720 ctx.abi.load_patch_vertices_in = load_patch_vertices_in;
3721 ctx.abi.store_tcs_outputs = store_tcs_output;
3722 ctx.tcs_vertices_per_patch = shaders[i]->info.tess.tcs_vertices_out;
3723 if (shader_count == 1)
3724 ctx.tcs_num_inputs = ctx.options->key.tcs.num_inputs;
3725 else
3726 ctx.tcs_num_inputs = util_last_bit64(shader_info->info.vs.ls_outputs_written);
3727 ctx.tcs_num_patches = get_tcs_num_patches(&ctx);
3728 } else if (shaders[i]->info.stage == MESA_SHADER_TESS_EVAL) {
3729 ctx.tes_primitive_mode = shaders[i]->info.tess.primitive_mode;
3730 ctx.abi.load_tess_varyings = load_tes_input;
3731 ctx.abi.load_tess_coord = load_tess_coord;
3732 ctx.abi.load_patch_vertices_in = load_patch_vertices_in;
3733 ctx.tcs_vertices_per_patch = shaders[i]->info.tess.tcs_vertices_out;
3734 ctx.tcs_num_patches = ctx.options->key.tes.num_patches;
3735 } else if (shaders[i]->info.stage == MESA_SHADER_VERTEX) {
3736 if (shader_info->info.vs.needs_instance_id) {
3737 if (ctx.options->key.vs.as_ls) {
3738 ctx.shader_info->vs.vgpr_comp_cnt =
3739 MAX2(2, ctx.shader_info->vs.vgpr_comp_cnt);
3740 } else {
3741 ctx.shader_info->vs.vgpr_comp_cnt =
3742 MAX2(1, ctx.shader_info->vs.vgpr_comp_cnt);
3743 }
3744 }
3745 ctx.abi.load_base_vertex = radv_load_base_vertex;
3746 } else if (shaders[i]->info.stage == MESA_SHADER_FRAGMENT) {
3747 shader_info->fs.can_discard = shaders[i]->info.fs.uses_discard;
3748 ctx.abi.lookup_interp_param = lookup_interp_param;
3749 ctx.abi.load_sample_position = load_sample_position;
3750 ctx.abi.load_sample_mask_in = load_sample_mask_in;
3751 ctx.abi.emit_kill = radv_emit_kill;
3752 }
3753
3754 if (i)
3755 ac_emit_barrier(&ctx.ac, ctx.stage);
3756
3757 nir_foreach_variable(variable, &shaders[i]->outputs)
3758 scan_shader_output_decl(&ctx, variable, shaders[i], shaders[i]->info.stage);
3759
3760 if (shaders[i]->info.stage == MESA_SHADER_GEOMETRY) {
3761 unsigned addclip = shaders[i]->info.clip_distance_array_size +
3762 shaders[i]->info.cull_distance_array_size > 4;
3763 ctx.gsvs_vertex_size = (util_bitcount64(ctx.output_mask) + addclip) * 16;
3764 ctx.max_gsvs_emit_size = ctx.gsvs_vertex_size *
3765 shaders[i]->info.gs.vertices_out;
3766 }
3767
3768 ac_setup_rings(&ctx);
3769
3770 LLVMBasicBlockRef merge_block;
3771 if (shader_count >= 2) {
3772 LLVMValueRef fn = LLVMGetBasicBlockParent(LLVMGetInsertBlock(ctx.ac.builder));
3773 LLVMBasicBlockRef then_block = LLVMAppendBasicBlockInContext(ctx.ac.context, fn, "");
3774 merge_block = LLVMAppendBasicBlockInContext(ctx.ac.context, fn, "");
3775
3776 LLVMValueRef count = ac_unpack_param(&ctx.ac, ctx.merged_wave_info, 8 * i, 8);
3777 LLVMValueRef thread_id = ac_get_thread_id(&ctx.ac);
3778 LLVMValueRef cond = LLVMBuildICmp(ctx.ac.builder, LLVMIntULT,
3779 thread_id, count, "");
3780 LLVMBuildCondBr(ctx.ac.builder, cond, then_block, merge_block);
3781
3782 LLVMPositionBuilderAtEnd(ctx.ac.builder, then_block);
3783 }
3784
3785 if (shaders[i]->info.stage == MESA_SHADER_FRAGMENT)
3786 handle_fs_inputs(&ctx, shaders[i]);
3787 else if(shaders[i]->info.stage == MESA_SHADER_VERTEX)
3788 handle_vs_inputs(&ctx, shaders[i]);
3789 else if(shader_count >= 2 && shaders[i]->info.stage == MESA_SHADER_GEOMETRY)
3790 prepare_gs_input_vgprs(&ctx);
3791
3792 ac_nir_translate(&ctx.ac, &ctx.abi, shaders[i]);
3793
3794 if (shader_count >= 2) {
3795 LLVMBuildBr(ctx.ac.builder, merge_block);
3796 LLVMPositionBuilderAtEnd(ctx.ac.builder, merge_block);
3797 }
3798
3799 if (shaders[i]->info.stage == MESA_SHADER_GEOMETRY) {
3800 shader_info->gs.gsvs_vertex_size = ctx.gsvs_vertex_size;
3801 shader_info->gs.max_gsvs_emit_size = ctx.max_gsvs_emit_size;
3802 } else if (shaders[i]->info.stage == MESA_SHADER_TESS_CTRL) {
3803 shader_info->tcs.num_patches = ctx.tcs_num_patches;
3804 shader_info->tcs.lds_size = calculate_tess_lds_size(&ctx);
3805 }
3806 }
3807
3808 LLVMBuildRetVoid(ctx.ac.builder);
3809
3810 if (options->dump_preoptir)
3811 ac_dump_module(ctx.ac.module);
3812
3813 ac_llvm_finalize_module(&ctx, ac_llvm->passmgr, options);
3814
3815 if (shader_count == 1)
3816 ac_nir_eliminate_const_vs_outputs(&ctx);
3817
3818 if (options->dump_shader) {
3819 ctx.shader_info->private_mem_vgprs =
3820 ac_count_scratch_private_memory(ctx.main_function);
3821 }
3822
3823 return ctx.ac.module;
3824 }
3825
3826 static void ac_diagnostic_handler(LLVMDiagnosticInfoRef di, void *context)
3827 {
3828 unsigned *retval = (unsigned *)context;
3829 LLVMDiagnosticSeverity severity = LLVMGetDiagInfoSeverity(di);
3830 char *description = LLVMGetDiagInfoDescription(di);
3831
3832 if (severity == LLVMDSError) {
3833 *retval = 1;
3834 fprintf(stderr, "LLVM triggered Diagnostic Handler: %s\n",
3835 description);
3836 }
3837
3838 LLVMDisposeMessage(description);
3839 }
3840
3841 static unsigned ac_llvm_compile(LLVMModuleRef M,
3842 struct ac_shader_binary *binary,
3843 struct ac_llvm_compiler *ac_llvm)
3844 {
3845 unsigned retval = 0;
3846 LLVMContextRef llvm_ctx;
3847
3848 /* Setup Diagnostic Handler*/
3849 llvm_ctx = LLVMGetModuleContext(M);
3850
3851 LLVMContextSetDiagnosticHandler(llvm_ctx, ac_diagnostic_handler,
3852 &retval);
3853
3854 /* Compile IR*/
3855 if (!radv_compile_to_binary(ac_llvm, M, binary))
3856 retval = 1;
3857 return retval;
3858 }
3859
3860 static void ac_compile_llvm_module(struct ac_llvm_compiler *ac_llvm,
3861 LLVMModuleRef llvm_module,
3862 struct ac_shader_binary *binary,
3863 struct ac_shader_config *config,
3864 struct radv_shader_variant_info *shader_info,
3865 gl_shader_stage stage,
3866 const struct radv_nir_compiler_options *options)
3867 {
3868 if (options->dump_shader)
3869 ac_dump_module(llvm_module);
3870
3871 memset(binary, 0, sizeof(*binary));
3872
3873 if (options->record_llvm_ir) {
3874 char *llvm_ir = LLVMPrintModuleToString(llvm_module);
3875 binary->llvm_ir_string = strdup(llvm_ir);
3876 LLVMDisposeMessage(llvm_ir);
3877 }
3878
3879 int v = ac_llvm_compile(llvm_module, binary, ac_llvm);
3880 if (v) {
3881 fprintf(stderr, "compile failed\n");
3882 }
3883
3884 if (options->dump_shader)
3885 fprintf(stderr, "disasm:\n%s\n", binary->disasm_string);
3886
3887 ac_shader_binary_read_config(binary, config, 0, options->supports_spill);
3888
3889 LLVMContextRef ctx = LLVMGetModuleContext(llvm_module);
3890 LLVMDisposeModule(llvm_module);
3891 LLVMContextDispose(ctx);
3892
3893 if (stage == MESA_SHADER_FRAGMENT) {
3894 shader_info->num_input_vgprs = 0;
3895 if (G_0286CC_PERSP_SAMPLE_ENA(config->spi_ps_input_addr))
3896 shader_info->num_input_vgprs += 2;
3897 if (G_0286CC_PERSP_CENTER_ENA(config->spi_ps_input_addr))
3898 shader_info->num_input_vgprs += 2;
3899 if (G_0286CC_PERSP_CENTROID_ENA(config->spi_ps_input_addr))
3900 shader_info->num_input_vgprs += 2;
3901 if (G_0286CC_PERSP_PULL_MODEL_ENA(config->spi_ps_input_addr))
3902 shader_info->num_input_vgprs += 3;
3903 if (G_0286CC_LINEAR_SAMPLE_ENA(config->spi_ps_input_addr))
3904 shader_info->num_input_vgprs += 2;
3905 if (G_0286CC_LINEAR_CENTER_ENA(config->spi_ps_input_addr))
3906 shader_info->num_input_vgprs += 2;
3907 if (G_0286CC_LINEAR_CENTROID_ENA(config->spi_ps_input_addr))
3908 shader_info->num_input_vgprs += 2;
3909 if (G_0286CC_LINE_STIPPLE_TEX_ENA(config->spi_ps_input_addr))
3910 shader_info->num_input_vgprs += 1;
3911 if (G_0286CC_POS_X_FLOAT_ENA(config->spi_ps_input_addr))
3912 shader_info->num_input_vgprs += 1;
3913 if (G_0286CC_POS_Y_FLOAT_ENA(config->spi_ps_input_addr))
3914 shader_info->num_input_vgprs += 1;
3915 if (G_0286CC_POS_Z_FLOAT_ENA(config->spi_ps_input_addr))
3916 shader_info->num_input_vgprs += 1;
3917 if (G_0286CC_POS_W_FLOAT_ENA(config->spi_ps_input_addr))
3918 shader_info->num_input_vgprs += 1;
3919 if (G_0286CC_FRONT_FACE_ENA(config->spi_ps_input_addr))
3920 shader_info->num_input_vgprs += 1;
3921 if (G_0286CC_ANCILLARY_ENA(config->spi_ps_input_addr))
3922 shader_info->num_input_vgprs += 1;
3923 if (G_0286CC_SAMPLE_COVERAGE_ENA(config->spi_ps_input_addr))
3924 shader_info->num_input_vgprs += 1;
3925 if (G_0286CC_POS_FIXED_PT_ENA(config->spi_ps_input_addr))
3926 shader_info->num_input_vgprs += 1;
3927 }
3928 config->num_vgprs = MAX2(config->num_vgprs, shader_info->num_input_vgprs);
3929
3930 /* +3 for scratch wave offset and VCC */
3931 config->num_sgprs = MAX2(config->num_sgprs,
3932 shader_info->num_input_sgprs + 3);
3933
3934 /* Enable 64-bit and 16-bit denormals, because there is no performance
3935 * cost.
3936 *
3937 * If denormals are enabled, all floating-point output modifiers are
3938 * ignored.
3939 *
3940 * Don't enable denormals for 32-bit floats, because:
3941 * - Floating-point output modifiers would be ignored by the hw.
3942 * - Some opcodes don't support denormals, such as v_mad_f32. We would
3943 * have to stop using those.
3944 * - SI & CI would be very slow.
3945 */
3946 config->float_mode |= V_00B028_FP_64_DENORMS;
3947 }
3948
3949 static void
3950 ac_fill_shader_info(struct radv_shader_variant_info *shader_info, struct nir_shader *nir, const struct radv_nir_compiler_options *options)
3951 {
3952 switch (nir->info.stage) {
3953 case MESA_SHADER_COMPUTE:
3954 for (int i = 0; i < 3; ++i)
3955 shader_info->cs.block_size[i] = nir->info.cs.local_size[i];
3956 break;
3957 case MESA_SHADER_FRAGMENT:
3958 shader_info->fs.early_fragment_test = nir->info.fs.early_fragment_tests;
3959 break;
3960 case MESA_SHADER_GEOMETRY:
3961 shader_info->gs.vertices_in = nir->info.gs.vertices_in;
3962 shader_info->gs.vertices_out = nir->info.gs.vertices_out;
3963 shader_info->gs.output_prim = nir->info.gs.output_primitive;
3964 shader_info->gs.invocations = nir->info.gs.invocations;
3965 break;
3966 case MESA_SHADER_TESS_EVAL:
3967 shader_info->tes.primitive_mode = nir->info.tess.primitive_mode;
3968 shader_info->tes.spacing = nir->info.tess.spacing;
3969 shader_info->tes.ccw = nir->info.tess.ccw;
3970 shader_info->tes.point_mode = nir->info.tess.point_mode;
3971 shader_info->tes.as_es = options->key.tes.as_es;
3972 break;
3973 case MESA_SHADER_TESS_CTRL:
3974 shader_info->tcs.tcs_vertices_out = nir->info.tess.tcs_vertices_out;
3975 break;
3976 case MESA_SHADER_VERTEX:
3977 shader_info->vs.as_es = options->key.vs.as_es;
3978 shader_info->vs.as_ls = options->key.vs.as_ls;
3979 /* in LS mode we need at least 1, invocation id needs 2, handled elsewhere */
3980 if (options->key.vs.as_ls)
3981 shader_info->vs.vgpr_comp_cnt = MAX2(1, shader_info->vs.vgpr_comp_cnt);
3982 break;
3983 default:
3984 break;
3985 }
3986 }
3987
3988 void
3989 radv_compile_nir_shader(struct ac_llvm_compiler *ac_llvm,
3990 struct ac_shader_binary *binary,
3991 struct ac_shader_config *config,
3992 struct radv_shader_variant_info *shader_info,
3993 struct nir_shader *const *nir,
3994 int nir_count,
3995 const struct radv_nir_compiler_options *options)
3996 {
3997
3998 LLVMModuleRef llvm_module;
3999
4000 llvm_module = ac_translate_nir_to_llvm(ac_llvm, nir, nir_count, shader_info,
4001 options);
4002
4003 ac_compile_llvm_module(ac_llvm, llvm_module, binary, config, shader_info,
4004 nir[0]->info.stage, options);
4005
4006 for (int i = 0; i < nir_count; ++i)
4007 ac_fill_shader_info(shader_info, nir[i], options);
4008
4009 /* Determine the ES type (VS or TES) for the GS on GFX9. */
4010 if (options->chip_class == GFX9) {
4011 if (nir_count == 2 &&
4012 nir[1]->info.stage == MESA_SHADER_GEOMETRY) {
4013 shader_info->gs.es_type = nir[0]->info.stage;
4014 }
4015 }
4016 }
4017
4018 static void
4019 ac_gs_copy_shader_emit(struct radv_shader_context *ctx)
4020 {
4021 LLVMValueRef vtx_offset =
4022 LLVMBuildMul(ctx->ac.builder, ctx->abi.vertex_id,
4023 LLVMConstInt(ctx->ac.i32, 4, false), "");
4024 LLVMValueRef stream_id;
4025
4026 /* Fetch the vertex stream ID. */
4027 if (ctx->shader_info->info.so.num_outputs) {
4028 stream_id =
4029 ac_unpack_param(&ctx->ac, ctx->streamout_config, 24, 2);
4030 } else {
4031 stream_id = ctx->ac.i32_0;
4032 }
4033
4034 LLVMBasicBlockRef end_bb;
4035 LLVMValueRef switch_inst;
4036
4037 end_bb = LLVMAppendBasicBlockInContext(ctx->ac.context,
4038 ctx->main_function, "end");
4039 switch_inst = LLVMBuildSwitch(ctx->ac.builder, stream_id, end_bb, 4);
4040
4041 for (unsigned stream = 0; stream < 4; stream++) {
4042 unsigned num_components =
4043 ctx->shader_info->info.gs.num_stream_output_components[stream];
4044 LLVMBasicBlockRef bb;
4045 unsigned offset;
4046
4047 if (!num_components)
4048 continue;
4049
4050 if (stream > 0 && !ctx->shader_info->info.so.num_outputs)
4051 continue;
4052
4053 bb = LLVMInsertBasicBlockInContext(ctx->ac.context, end_bb, "out");
4054 LLVMAddCase(switch_inst, LLVMConstInt(ctx->ac.i32, stream, 0), bb);
4055 LLVMPositionBuilderAtEnd(ctx->ac.builder, bb);
4056
4057 offset = 0;
4058 for (unsigned i = 0; i < AC_LLVM_MAX_OUTPUTS; ++i) {
4059 unsigned output_usage_mask =
4060 ctx->shader_info->info.gs.output_usage_mask[i];
4061 unsigned output_stream =
4062 ctx->shader_info->info.gs.output_streams[i];
4063 int length = util_last_bit(output_usage_mask);
4064
4065 if (!(ctx->output_mask & (1ull << i)) ||
4066 output_stream != stream)
4067 continue;
4068
4069 for (unsigned j = 0; j < length; j++) {
4070 LLVMValueRef value, soffset;
4071
4072 if (!(output_usage_mask & (1 << j)))
4073 continue;
4074
4075 soffset = LLVMConstInt(ctx->ac.i32,
4076 offset *
4077 ctx->gs_max_out_vertices * 16 * 4, false);
4078
4079 offset++;
4080
4081 value = ac_build_buffer_load(&ctx->ac,
4082 ctx->gsvs_ring[0],
4083 1, ctx->ac.i32_0,
4084 vtx_offset, soffset,
4085 0, 1, 1, true, false);
4086
4087 LLVMTypeRef type = LLVMGetAllocatedType(ctx->abi.outputs[ac_llvm_reg_index_soa(i, j)]);
4088 if (ac_get_type_size(type) == 2) {
4089 value = LLVMBuildBitCast(ctx->ac.builder, value, ctx->ac.i32, "");
4090 value = LLVMBuildTrunc(ctx->ac.builder, value, ctx->ac.i16, "");
4091 }
4092
4093 LLVMBuildStore(ctx->ac.builder,
4094 ac_to_float(&ctx->ac, value), ctx->abi.outputs[ac_llvm_reg_index_soa(i, j)]);
4095 }
4096 }
4097
4098 if (ctx->shader_info->info.so.num_outputs)
4099 radv_emit_streamout(ctx, stream);
4100
4101 if (stream == 0) {
4102 handle_vs_outputs_post(ctx, false, false,
4103 &ctx->shader_info->vs.outinfo);
4104 }
4105
4106 LLVMBuildBr(ctx->ac.builder, end_bb);
4107 }
4108
4109 LLVMPositionBuilderAtEnd(ctx->ac.builder, end_bb);
4110 }
4111
4112 void
4113 radv_compile_gs_copy_shader(struct ac_llvm_compiler *ac_llvm,
4114 struct nir_shader *geom_shader,
4115 struct ac_shader_binary *binary,
4116 struct ac_shader_config *config,
4117 struct radv_shader_variant_info *shader_info,
4118 const struct radv_nir_compiler_options *options)
4119 {
4120 struct radv_shader_context ctx = {0};
4121 ctx.options = options;
4122 ctx.shader_info = shader_info;
4123
4124 ac_llvm_context_init(&ctx.ac, options->chip_class, options->family);
4125 ctx.context = ctx.ac.context;
4126 ctx.ac.module = ac_create_module(ac_llvm->tm, ctx.context);
4127
4128 ctx.is_gs_copy_shader = true;
4129
4130 enum ac_float_mode float_mode =
4131 options->unsafe_math ? AC_FLOAT_MODE_UNSAFE_FP_MATH :
4132 AC_FLOAT_MODE_DEFAULT;
4133
4134 ctx.ac.builder = ac_create_builder(ctx.context, float_mode);
4135 ctx.stage = MESA_SHADER_VERTEX;
4136
4137 radv_nir_shader_info_pass(geom_shader, options, &shader_info->info);
4138
4139 create_function(&ctx, MESA_SHADER_VERTEX, false, MESA_SHADER_VERTEX);
4140
4141 ctx.gs_max_out_vertices = geom_shader->info.gs.vertices_out;
4142 ac_setup_rings(&ctx);
4143
4144 nir_foreach_variable(variable, &geom_shader->outputs) {
4145 scan_shader_output_decl(&ctx, variable, geom_shader, MESA_SHADER_VERTEX);
4146 ac_handle_shader_output_decl(&ctx.ac, &ctx.abi, geom_shader,
4147 variable, MESA_SHADER_VERTEX);
4148 }
4149
4150 ac_gs_copy_shader_emit(&ctx);
4151
4152 LLVMBuildRetVoid(ctx.ac.builder);
4153
4154 ac_llvm_finalize_module(&ctx, ac_llvm->passmgr, options);
4155
4156 ac_compile_llvm_module(ac_llvm, ctx.ac.module, binary, config, shader_info,
4157 MESA_SHADER_VERTEX, options);
4158 }