1 /*
2 * Copyright 2012 Advanced Micro Devices, Inc.
3 *
4 * SPDX-License-Identifier: MIT
5 */
6
7 #if AMD_LLVM_AVAILABLE
8 #include "ac_llvm_util.h"
9 #endif
10
11 #include "ac_nir.h"
12 #include "ac_shader_util.h"
13 #include "compiler/nir/nir_serialize.h"
14 #include "nir/tgsi_to_nir.h"
15 #include "si_build_pm4.h"
16 #include "sid.h"
17 #include "util/crc32.h"
18 #include "util/disk_cache.h"
19 #include "util/hash_table.h"
20 #include "util/mesa-sha1.h"
21 #include "util/u_async_debug.h"
22 #include "util/u_math.h"
23 #include "util/u_memory.h"
24 #include "util/u_prim.h"
25 #include "tgsi/tgsi_from_mesa.h"
26
27 static void si_update_tess_in_out_patch_vertices(struct si_context *sctx);
28
si_determine_wave_size(struct si_screen * sscreen,struct si_shader * shader)29 unsigned si_determine_wave_size(struct si_screen *sscreen, struct si_shader *shader)
30 {
31 struct si_shader_info *info = &shader->selector->info;
32 gl_shader_stage stage = shader->selector->stage;
33
34 struct si_shader_selector *prev_sel = NULL;
35 if (stage == MESA_SHADER_TESS_CTRL)
36 prev_sel = shader->key.ge.part.tcs.ls;
37 else if (stage == MESA_SHADER_GEOMETRY)
38 prev_sel = shader->key.ge.part.gs.es;
39
40 if (sscreen->info.gfx_level < GFX10)
41 return 64;
42
43 /* Legacy GS only supports Wave64. */
44 if ((stage == MESA_SHADER_VERTEX && shader->key.ge.as_es && !shader->key.ge.as_ngg) ||
45 (stage == MESA_SHADER_TESS_EVAL && shader->key.ge.as_es && !shader->key.ge.as_ngg) ||
46 (stage == MESA_SHADER_GEOMETRY && !shader->key.ge.as_ngg))
47 return 64;
48
49 /* For KHR_shader_subgroup which require a constant subgroup size known by user. */
50 if (info->base.subgroup_size == SUBGROUP_SIZE_API_CONSTANT ||
51 (prev_sel && prev_sel->info.base.subgroup_size == SUBGROUP_SIZE_API_CONSTANT))
52 return 64;
53
54 /* Workgroup sizes that are not divisible by 64 use Wave32. */
55 if (stage == MESA_SHADER_COMPUTE && !info->base.workgroup_size_variable &&
56 (info->base.workgroup_size[0] *
57 info->base.workgroup_size[1] *
58 info->base.workgroup_size[2]) % 64 != 0)
59 return 32;
60
61 /* AMD_DEBUG wave flags override everything else. */
62 if (sscreen->debug_flags &
63 (stage == MESA_SHADER_COMPUTE ? DBG(W32_CS) :
64 stage == MESA_SHADER_FRAGMENT ? DBG(W32_PS) : DBG(W32_GE)))
65 return 32;
66
67 if (sscreen->debug_flags &
68 (stage == MESA_SHADER_COMPUTE ? DBG(W64_CS) :
69 stage == MESA_SHADER_FRAGMENT ? DBG(W64_PS) : DBG(W64_GE)))
70 return 64;
71
72 /* Shader profiles. */
73 if (info->options & SI_PROFILE_WAVE32)
74 return 32;
75
76 if (info->options & SI_PROFILE_GFX10_WAVE64 &&
77 (sscreen->info.gfx_level == GFX10 || sscreen->info.gfx_level == GFX10_3))
78 return 64;
79
80 /* Gfx10: Pixel shaders without interp instructions don't suffer from reduced interpolation
81 * performance in Wave32, so use Wave32. This helps Piano and Voloplosion.
82 *
83 * Gfx11: Prefer Wave64 to take advantage of doubled VALU performance.
84 */
85 if (sscreen->info.gfx_level < GFX11 && stage == MESA_SHADER_FRAGMENT && !info->num_inputs)
86 return 32;
87
88 /* Gfx10: There are a few very rare cases where VS is better with Wave32, and there are no
89 * known cases where Wave64 is better.
90 *
91 * Wave32 is disabled for GFX10 when culling is active as a workaround for #6457. I don't
92 * know why this helps.
93 *
94 * Gfx11: Prefer Wave64 because it's slightly better than Wave32.
95 */
96 if (stage <= MESA_SHADER_GEOMETRY &&
97 (sscreen->info.gfx_level == GFX10 || sscreen->info.gfx_level == GFX10_3) &&
98 !(sscreen->info.gfx_level == GFX10 && si_shader_culling_enabled(shader)))
99 return 32;
100
101 /* Divergent loops in Wave64 can end up having too many iterations in one half of the wave
102 * while the other half is idling but occupying VGPRs, preventing other waves from launching.
103 * Wave32 eliminates the idling half to allow the next wave to start.
104 *
105 * Gfx11: Wave32 continues to be faster with divergent loops despite worse VALU performance.
106 */
107 if (info->has_divergent_loop ||
108 /* Merged shader has to use same wave size for two shader stages. */
109 (prev_sel && prev_sel->info.has_divergent_loop))
110 return 32;
111
112 return 64;
113 }
114
si_shader_uses_bindless_samplers(struct si_shader_selector * selector)115 static bool si_shader_uses_bindless_samplers(struct si_shader_selector *selector)
116 {
117 return selector ? selector->info.uses_bindless_samplers : false;
118 }
119
si_shader_uses_bindless_images(struct si_shader_selector * selector)120 static bool si_shader_uses_bindless_images(struct si_shader_selector *selector)
121 {
122 return selector ? selector->info.uses_bindless_images : false;
123 }
124
125 /* SHADER_CACHE */
126
127 /**
128 * Return the IR key for the shader cache.
129 */
si_get_ir_cache_key(struct si_shader_selector * sel,bool ngg,bool es,unsigned wave_size,unsigned char ir_sha1_cache_key[20])130 void si_get_ir_cache_key(struct si_shader_selector *sel, bool ngg, bool es,
131 unsigned wave_size, unsigned char ir_sha1_cache_key[20])
132 {
133 struct blob blob = {};
134 unsigned ir_size;
135 void *ir_binary;
136
137 if (sel->nir_binary) {
138 ir_binary = sel->nir_binary;
139 ir_size = sel->nir_size;
140 } else {
141 assert(sel->nir);
142
143 blob_init(&blob);
144 /* Keep debug info if NIR debug prints are in use. */
145 nir_serialize(&blob, sel->nir, NIR_DEBUG(PRINT) == 0);
146 ir_binary = blob.data;
147 ir_size = blob.size;
148 }
149
150 /* These settings affect the compilation, but they are not derived
151 * from the input shader IR.
152 */
153 unsigned shader_variant_flags = 0;
154
155 if (ngg)
156 shader_variant_flags |= 1 << 0;
157 /* bit gap */
158 if (wave_size == 32)
159 shader_variant_flags |= 1 << 2;
160 /* bit gap */
161 /* use_ngg_culling disables NGG passthrough for non-culling shaders to reduce context
162 * rolls, which can be changed with AMD_DEBUG=nonggc or AMD_DEBUG=nggc.
163 */
164 if (sel->screen->use_ngg_culling)
165 shader_variant_flags |= 1 << 4;
166 if (sel->screen->record_llvm_ir)
167 shader_variant_flags |= 1 << 5;
168 if (sel->screen->info.has_image_opcodes)
169 shader_variant_flags |= 1 << 6;
170 if (sel->screen->options.no_infinite_interp)
171 shader_variant_flags |= 1 << 7;
172 if (sel->screen->options.clamp_div_by_zero)
173 shader_variant_flags |= 1 << 8;
174 if ((sel->stage == MESA_SHADER_VERTEX ||
175 sel->stage == MESA_SHADER_TESS_EVAL ||
176 sel->stage == MESA_SHADER_GEOMETRY) &&
177 !es &&
178 sel->screen->options.vrs2x2)
179 shader_variant_flags |= 1 << 10;
180 if (sel->screen->options.inline_uniforms)
181 shader_variant_flags |= 1 << 11;
182 if (sel->screen->options.clear_lds)
183 shader_variant_flags |= 1 << 12;
184
185 struct mesa_sha1 ctx;
186 _mesa_sha1_init(&ctx);
187 _mesa_sha1_update(&ctx, &shader_variant_flags, 4);
188 _mesa_sha1_update(&ctx, ir_binary, ir_size);
189 _mesa_sha1_final(&ctx, ir_sha1_cache_key);
190
191 if (ir_binary == blob.data)
192 blob_finish(&blob);
193 }
194
195 /** Copy "data" to "ptr" and return the next dword following copied data. */
write_data(uint32_t * ptr,const void * data,unsigned size)196 static uint32_t *write_data(uint32_t *ptr, const void *data, unsigned size)
197 {
198 /* data may be NULL if size == 0 */
199 if (size)
200 memcpy(ptr, data, size);
201 ptr += DIV_ROUND_UP(size, 4);
202 return ptr;
203 }
204
205 /** Read data from "ptr". Return the next dword following the data. */
read_data(uint32_t * ptr,void * data,unsigned size)206 static uint32_t *read_data(uint32_t *ptr, void *data, unsigned size)
207 {
208 memcpy(data, ptr, size);
209 ptr += DIV_ROUND_UP(size, 4);
210 return ptr;
211 }
212
213 /**
214 * Write the size as uint followed by the data. Return the next dword
215 * following the copied data.
216 */
write_chunk(uint32_t * ptr,const void * data,unsigned size)217 static uint32_t *write_chunk(uint32_t *ptr, const void *data, unsigned size)
218 {
219 *ptr++ = size;
220 return write_data(ptr, data, size);
221 }
222
223 /**
224 * Read the size as uint followed by the data. Return both via parameters.
225 * Return the next dword following the data.
226 */
read_chunk(uint32_t * ptr,void ** data,unsigned * size)227 static uint32_t *read_chunk(uint32_t *ptr, void **data, unsigned *size)
228 {
229 *size = *ptr++;
230 assert(*data == NULL);
231 if (!*size)
232 return ptr;
233 *data = malloc(*size);
234 return read_data(ptr, *data, *size);
235 }
236
237 struct si_shader_blob_head {
238 uint32_t size;
239 uint32_t type;
240 uint32_t crc32;
241 };
242
243 /**
244 * Return the shader binary in a buffer.
245 */
si_get_shader_binary(struct si_shader * shader)246 static uint32_t *si_get_shader_binary(struct si_shader *shader)
247 {
248 /* There is always a size of data followed by the data itself. */
249 unsigned llvm_ir_size =
250 shader->binary.llvm_ir_string ? strlen(shader->binary.llvm_ir_string) + 1 : 0;
251
252 /* Refuse to allocate overly large buffers and guard against integer
253 * overflow. */
254 if (shader->binary.code_size > UINT_MAX / 4 || llvm_ir_size > UINT_MAX / 4 ||
255 shader->binary.num_symbols > UINT_MAX / 32)
256 return NULL;
257
258 unsigned size = sizeof(struct si_shader_blob_head) +
259 align(sizeof(shader->config), 4) +
260 align(sizeof(shader->info), 4) +
261 4 + 4 + align(shader->binary.code_size, 4) +
262 4 + shader->binary.num_symbols * 8 +
263 4 + align(llvm_ir_size, 4) +
264 4 + align(shader->binary.disasm_size, 4);
265 uint32_t *buffer = (uint32_t*)CALLOC(1, size);
266 if (!buffer)
267 return NULL;
268
269 struct si_shader_blob_head *head = (struct si_shader_blob_head *)buffer;
270 head->type = shader->binary.type;
271 head->size = size;
272
273 uint32_t *data = buffer + sizeof(*head) / 4;
274 uint32_t *ptr = data;
275
276 ptr = write_data(ptr, &shader->config, sizeof(shader->config));
277 ptr = write_data(ptr, &shader->info, sizeof(shader->info));
278 ptr = write_data(ptr, &shader->binary.exec_size, 4);
279 ptr = write_chunk(ptr, shader->binary.code_buffer, shader->binary.code_size);
280 ptr = write_chunk(ptr, shader->binary.symbols, shader->binary.num_symbols * 8);
281 ptr = write_chunk(ptr, shader->binary.llvm_ir_string, llvm_ir_size);
282 ptr = write_chunk(ptr, shader->binary.disasm_string, shader->binary.disasm_size);
283 assert((char *)ptr - (char *)buffer == (ptrdiff_t)size);
284
285 /* Compute CRC32. */
286 head->crc32 = util_hash_crc32(data, size - sizeof(*head));
287
288 return buffer;
289 }
290
si_load_shader_binary(struct si_shader * shader,void * binary)291 static bool si_load_shader_binary(struct si_shader *shader, void *binary)
292 {
293 struct si_shader_blob_head *head = (struct si_shader_blob_head *)binary;
294 unsigned chunk_size;
295 unsigned code_size;
296
297 uint32_t *ptr = (uint32_t *)binary + sizeof(*head) / 4;
298 if (util_hash_crc32(ptr, head->size - sizeof(*head)) != head->crc32) {
299 fprintf(stderr, "radeonsi: binary shader has invalid CRC32\n");
300 return false;
301 }
302
303 shader->binary.type = (enum si_shader_binary_type)head->type;
304 ptr = read_data(ptr, &shader->config, sizeof(shader->config));
305 ptr = read_data(ptr, &shader->info, sizeof(shader->info));
306 ptr = read_data(ptr, &shader->binary.exec_size, 4);
307 ptr = read_chunk(ptr, (void **)&shader->binary.code_buffer, &code_size);
308 shader->binary.code_size = code_size;
309 ptr = read_chunk(ptr, (void **)&shader->binary.symbols, &chunk_size);
310 shader->binary.num_symbols = chunk_size / 8;
311 ptr = read_chunk(ptr, (void **)&shader->binary.llvm_ir_string, &chunk_size);
312 ptr = read_chunk(ptr, (void **)&shader->binary.disasm_string, &chunk_size);
313 shader->binary.disasm_size = chunk_size;
314
315 if (!shader->is_gs_copy_shader &&
316 shader->selector->stage == MESA_SHADER_GEOMETRY && !shader->key.ge.as_ngg) {
317 shader->gs_copy_shader = CALLOC_STRUCT(si_shader);
318 if (!shader->gs_copy_shader)
319 return false;
320
321 shader->gs_copy_shader->is_gs_copy_shader = true;
322
323 if (!si_load_shader_binary(shader->gs_copy_shader, (uint8_t*)binary + head->size)) {
324 FREE(shader->gs_copy_shader);
325 shader->gs_copy_shader = NULL;
326 return false;
327 }
328
329 util_queue_fence_init(&shader->gs_copy_shader->ready);
330 shader->gs_copy_shader->selector = shader->selector;
331 shader->gs_copy_shader->is_gs_copy_shader = true;
332 shader->gs_copy_shader->wave_size =
333 si_determine_wave_size(shader->selector->screen, shader->gs_copy_shader);
334
335 si_shader_binary_upload(shader->selector->screen, shader->gs_copy_shader, 0);
336 }
337
338 return true;
339 }
340
341 /**
342 * Insert a shader into the cache. It's assumed the shader is not in the cache.
343 * Use si_shader_cache_load_shader before calling this.
344 */
si_shader_cache_insert_shader(struct si_screen * sscreen,unsigned char ir_sha1_cache_key[20],struct si_shader * shader,bool insert_into_disk_cache)345 void si_shader_cache_insert_shader(struct si_screen *sscreen, unsigned char ir_sha1_cache_key[20],
346 struct si_shader *shader, bool insert_into_disk_cache)
347 {
348 uint32_t *hw_binary;
349 struct hash_entry *entry;
350 uint8_t key[CACHE_KEY_SIZE];
351 bool memory_cache_full = sscreen->shader_cache_size >= sscreen->shader_cache_max_size;
352
353 if (!insert_into_disk_cache && memory_cache_full)
354 return;
355
356 entry = _mesa_hash_table_search(sscreen->shader_cache, ir_sha1_cache_key);
357 if (entry)
358 return; /* already added */
359
360 hw_binary = si_get_shader_binary(shader);
361 if (!hw_binary)
362 return;
363
364 unsigned size = *hw_binary;
365
366 if (shader->selector->stage == MESA_SHADER_GEOMETRY && !shader->key.ge.as_ngg) {
367 uint32_t *gs_copy_binary = si_get_shader_binary(shader->gs_copy_shader);
368 if (!gs_copy_binary) {
369 FREE(hw_binary);
370 return;
371 }
372
373 /* Combine both binaries. */
374 size += *gs_copy_binary;
375 uint32_t *combined_binary = (uint32_t*)MALLOC(size);
376 if (!combined_binary) {
377 FREE(hw_binary);
378 FREE(gs_copy_binary);
379 return;
380 }
381
382 memcpy(combined_binary, hw_binary, *hw_binary);
383 memcpy(combined_binary + *hw_binary / 4, gs_copy_binary, *gs_copy_binary);
384 FREE(hw_binary);
385 FREE(gs_copy_binary);
386 hw_binary = combined_binary;
387 }
388
389 if (!memory_cache_full) {
390 if (_mesa_hash_table_insert(sscreen->shader_cache,
391 mem_dup(ir_sha1_cache_key, 20),
392 hw_binary) == NULL) {
393 FREE(hw_binary);
394 return;
395 }
396
397 sscreen->shader_cache_size += size;
398 }
399
400 if (sscreen->disk_shader_cache && insert_into_disk_cache) {
401 disk_cache_compute_key(sscreen->disk_shader_cache, ir_sha1_cache_key, 20, key);
402 disk_cache_put(sscreen->disk_shader_cache, key, hw_binary, size, NULL);
403 }
404
405 if (memory_cache_full)
406 FREE(hw_binary);
407 }
408
si_shader_cache_load_shader(struct si_screen * sscreen,unsigned char ir_sha1_cache_key[20],struct si_shader * shader)409 bool si_shader_cache_load_shader(struct si_screen *sscreen, unsigned char ir_sha1_cache_key[20],
410 struct si_shader *shader)
411 {
412 struct hash_entry *entry = _mesa_hash_table_search(sscreen->shader_cache, ir_sha1_cache_key);
413
414 if (entry) {
415 if (si_load_shader_binary(shader, entry->data)) {
416 p_atomic_inc(&sscreen->num_memory_shader_cache_hits);
417 return true;
418 }
419 }
420 p_atomic_inc(&sscreen->num_memory_shader_cache_misses);
421
422 if (!sscreen->disk_shader_cache)
423 return false;
424
425 unsigned char sha1[CACHE_KEY_SIZE];
426 disk_cache_compute_key(sscreen->disk_shader_cache, ir_sha1_cache_key, 20, sha1);
427
428 size_t total_size;
429 uint32_t *buffer = (uint32_t*)disk_cache_get(sscreen->disk_shader_cache, sha1, &total_size);
430 if (buffer) {
431 unsigned size = *buffer;
432 unsigned gs_copy_binary_size = 0;
433
434 /* The GS copy shader binary is after the GS binary. */
435 if (shader->selector->stage == MESA_SHADER_GEOMETRY && !shader->key.ge.as_ngg)
436 gs_copy_binary_size = buffer[size / 4];
437
438 if (total_size >= sizeof(uint32_t) && size + gs_copy_binary_size == total_size) {
439 if (si_load_shader_binary(shader, buffer)) {
440 free(buffer);
441 si_shader_cache_insert_shader(sscreen, ir_sha1_cache_key, shader, false);
442 p_atomic_inc(&sscreen->num_disk_shader_cache_hits);
443 return true;
444 }
445 } else {
446 /* Something has gone wrong discard the item from the cache and
447 * rebuild/link from source.
448 */
449 assert(!"Invalid radeonsi shader disk cache item!");
450 disk_cache_remove(sscreen->disk_shader_cache, sha1);
451 }
452 }
453
454 free(buffer);
455 p_atomic_inc(&sscreen->num_disk_shader_cache_misses);
456 return false;
457 }
458
si_shader_cache_key_hash(const void * key)459 static uint32_t si_shader_cache_key_hash(const void *key)
460 {
461 /* Take the first dword of SHA1. */
462 return *(uint32_t *)key;
463 }
464
si_shader_cache_key_equals(const void * a,const void * b)465 static bool si_shader_cache_key_equals(const void *a, const void *b)
466 {
467 /* Compare SHA1s. */
468 return memcmp(a, b, 20) == 0;
469 }
470
si_destroy_shader_cache_entry(struct hash_entry * entry)471 static void si_destroy_shader_cache_entry(struct hash_entry *entry)
472 {
473 FREE((void *)entry->key);
474 FREE(entry->data);
475 }
476
si_init_shader_cache(struct si_screen * sscreen)477 bool si_init_shader_cache(struct si_screen *sscreen)
478 {
479 (void)simple_mtx_init(&sscreen->shader_cache_mutex, mtx_plain);
480 sscreen->shader_cache =
481 _mesa_hash_table_create(NULL, si_shader_cache_key_hash, si_shader_cache_key_equals);
482 sscreen->shader_cache_size = 0;
483 /* Maximum size: 64MB on 32 bits, 1GB else */
484 sscreen->shader_cache_max_size = ((sizeof(void *) == 4) ? 64 : 1024) * 1024 * 1024;
485
486 return sscreen->shader_cache != NULL;
487 }
488
si_destroy_shader_cache(struct si_screen * sscreen)489 void si_destroy_shader_cache(struct si_screen *sscreen)
490 {
491 if (sscreen->shader_cache)
492 _mesa_hash_table_destroy(sscreen->shader_cache, si_destroy_shader_cache_entry);
493 simple_mtx_destroy(&sscreen->shader_cache_mutex);
494 }
495
496 /* SHADER STATES */
497
si_shader_encode_vgprs(struct si_shader * shader)498 unsigned si_shader_encode_vgprs(struct si_shader *shader)
499 {
500 assert(shader->selector->screen->info.gfx_level >= GFX10 || shader->wave_size == 64);
501 return shader->config.num_vgprs / (shader->wave_size == 32 ? 8 : 4) - 1;
502 }
503
si_shader_encode_sgprs(struct si_shader * shader)504 unsigned si_shader_encode_sgprs(struct si_shader *shader)
505 {
506 if (shader->selector->screen->info.gfx_level >= GFX10)
507 return 0; /* Gfx10+ don't have the SGPRS field and always allocate 128 SGPRs. */
508
509 return shader->config.num_sgprs / 8 - 1;
510 }
511
si_shader_mem_ordered(struct si_shader * shader)512 bool si_shader_mem_ordered(struct si_shader *shader)
513 {
514 struct si_screen *sscreen = shader->selector->screen;
515
516 if (sscreen->info.gfx_level < GFX10 || sscreen->info.gfx_level >= GFX12)
517 return false;
518
519 /* Return true if both types of VMEM that return something are used. */
520 return shader->info.uses_vmem_sampler_or_bvh &&
521 (shader->info.uses_vmem_load_other ||
522 shader->config.scratch_bytes_per_wave);
523 }
524
si_set_tesseval_regs(struct si_screen * sscreen,const struct si_shader_selector * tes,struct si_shader * shader)525 static void si_set_tesseval_regs(struct si_screen *sscreen, const struct si_shader_selector *tes,
526 struct si_shader *shader)
527 {
528 const struct si_shader_info *info = &tes->info;
529 enum tess_primitive_mode tes_prim_mode = info->base.tess._primitive_mode;
530 unsigned tes_spacing = info->base.tess.spacing;
531 bool tes_vertex_order_cw = !info->base.tess.ccw;
532 bool tes_point_mode = info->base.tess.point_mode;
533 unsigned type, partitioning, topology, distribution_mode;
534
535 switch (tes_prim_mode) {
536 case TESS_PRIMITIVE_ISOLINES:
537 type = V_028B6C_TESS_ISOLINE;
538 break;
539 case TESS_PRIMITIVE_TRIANGLES:
540 type = V_028B6C_TESS_TRIANGLE;
541 break;
542 case TESS_PRIMITIVE_QUADS:
543 type = V_028B6C_TESS_QUAD;
544 break;
545 default:
546 assert(0);
547 return;
548 }
549
550 switch (tes_spacing) {
551 case TESS_SPACING_FRACTIONAL_ODD:
552 partitioning = V_028B6C_PART_FRAC_ODD;
553 break;
554 case TESS_SPACING_FRACTIONAL_EVEN:
555 partitioning = V_028B6C_PART_FRAC_EVEN;
556 break;
557 case TESS_SPACING_EQUAL:
558 partitioning = V_028B6C_PART_INTEGER;
559 break;
560 default:
561 assert(0);
562 return;
563 }
564
565 if (tes_point_mode)
566 topology = V_028B6C_OUTPUT_POINT;
567 else if (tes_prim_mode == TESS_PRIMITIVE_ISOLINES)
568 topology = V_028B6C_OUTPUT_LINE;
569 else if (tes_vertex_order_cw)
570 /* for some reason, this must be the other way around */
571 topology = V_028B6C_OUTPUT_TRIANGLE_CCW;
572 else
573 topology = V_028B6C_OUTPUT_TRIANGLE_CW;
574
575 if (sscreen->info.has_distributed_tess) {
576 if (sscreen->info.family == CHIP_FIJI || sscreen->info.family >= CHIP_POLARIS10)
577 distribution_mode = V_028B6C_TRAPEZOIDS;
578 else
579 distribution_mode = V_028B6C_DONUTS;
580 } else
581 distribution_mode = V_028B6C_NO_DIST;
582
583 shader->vgt_tf_param = S_028B6C_TYPE(type) | S_028B6C_PARTITIONING(partitioning) |
584 S_028B6C_TOPOLOGY(topology) |
585 S_028B6C_DISTRIBUTION_MODE(distribution_mode);
586
587 if (sscreen->info.gfx_level >= GFX12)
588 shader->vgt_tf_param |= S_028AA4_TEMPORAL(gfx12_load_last_use_discard);
589 }
590
591 /* Polaris needs different VTX_REUSE_DEPTH settings depending on
592 * whether the "fractional odd" tessellation spacing is used.
593 *
594 * Possible VGT configurations and which state should set the register:
595 *
596 * Reg set in | VGT shader configuration | Value
597 * ------------------------------------------------------
598 * VS as VS | VS | 30
599 * VS as ES | ES -> GS -> VS | 30
600 * TES as VS | LS -> HS -> VS | 14 or 30
601 * TES as ES | LS -> HS -> ES -> GS -> VS | 14 or 30
602 */
polaris_set_vgt_vertex_reuse(struct si_screen * sscreen,struct si_shader_selector * sel,struct si_shader * shader)603 static void polaris_set_vgt_vertex_reuse(struct si_screen *sscreen, struct si_shader_selector *sel,
604 struct si_shader *shader)
605 {
606 if (sscreen->info.family < CHIP_POLARIS10 || sscreen->info.gfx_level >= GFX10)
607 return;
608
609 /* VS as VS, or VS as ES: */
610 if ((sel->stage == MESA_SHADER_VERTEX &&
611 (!shader->key.ge.as_ls && !shader->is_gs_copy_shader)) ||
612 /* TES as VS, or TES as ES: */
613 sel->stage == MESA_SHADER_TESS_EVAL) {
614 unsigned vtx_reuse_depth = 30;
615
616 if (sel->stage == MESA_SHADER_TESS_EVAL &&
617 sel->info.base.tess.spacing == TESS_SPACING_FRACTIONAL_ODD)
618 vtx_reuse_depth = 14;
619
620 shader->vgt_vertex_reuse_block_cntl = vtx_reuse_depth;
621 }
622 }
623
624 static struct si_pm4_state *
si_get_shader_pm4_state(struct si_shader * shader,void (* emit_func)(struct si_context * ctx,unsigned index))625 si_get_shader_pm4_state(struct si_shader *shader,
626 void (*emit_func)(struct si_context *ctx, unsigned index))
627 {
628 si_pm4_clear_state(&shader->pm4, shader->selector->screen, false);
629 shader->pm4.atom.emit = emit_func;
630 return &shader->pm4;
631 }
632
si_get_num_vs_user_sgprs(struct si_shader * shader,unsigned num_always_on_user_sgprs)633 static unsigned si_get_num_vs_user_sgprs(struct si_shader *shader,
634 unsigned num_always_on_user_sgprs)
635 {
636 struct si_shader_selector *vs =
637 shader->previous_stage_sel ? shader->previous_stage_sel : shader->selector;
638 unsigned num_vbos_in_user_sgprs = vs->info.num_vbos_in_user_sgprs;
639
640 /* 1 SGPR is reserved for the vertex buffer pointer. */
641 assert(num_always_on_user_sgprs <= SI_SGPR_VS_VB_DESCRIPTOR_FIRST - 1);
642
643 if (num_vbos_in_user_sgprs)
644 return SI_SGPR_VS_VB_DESCRIPTOR_FIRST + num_vbos_in_user_sgprs * 4;
645
646 /* Add the pointer to VBO descriptors. */
647 return num_always_on_user_sgprs + 1;
648 }
649
650 /* Return VGPR_COMP_CNT for the API vertex shader. This can be hw LS, LSHS, ES, ESGS, VS. */
si_get_vs_vgpr_comp_cnt(struct si_screen * sscreen,struct si_shader * shader,bool legacy_vs_prim_id)651 static unsigned si_get_vs_vgpr_comp_cnt(struct si_screen *sscreen, struct si_shader *shader,
652 bool legacy_vs_prim_id)
653 {
654 assert(shader->selector->stage == MESA_SHADER_VERTEX ||
655 (shader->previous_stage_sel && shader->previous_stage_sel->stage == MESA_SHADER_VERTEX));
656
657 /* GFX6-9 LS (VertexID, RelAutoIndex, InstanceID / StepRate0, InstanceID)
658 * GFX6-9 ES,VS (VertexID, InstanceID / StepRate0, VSPrimID, InstanceID)
659 * GFX10-11 LS (VertexID, RelAutoIndex, UserVGPR1, UserVGPR2 or InstanceID)
660 * GFX10-11 ES,VS (VertexID, UserVGPR1, UserVGPR2 or VSPrimID, UserVGPR3 or InstanceID)
661 * GFX12 LS,ES (VertexID, InstanceID)
662 */
663 bool is_ls = shader->selector->stage == MESA_SHADER_TESS_CTRL || shader->key.ge.as_ls;
664 unsigned max = 0;
665
666 if (shader->info.uses_instanceid) {
667 if (sscreen->info.gfx_level >= GFX12)
668 max = MAX2(max, 1);
669 else if (sscreen->info.gfx_level >= GFX10)
670 max = MAX2(max, 3);
671 else if (is_ls)
672 max = MAX2(max, 2); /* use (InstanceID / StepRate0) because StepRate0 == 1 */
673 else
674 max = MAX2(max, 1); /* use (InstanceID / StepRate0) because StepRate0 == 1 */
675 }
676
677 if (legacy_vs_prim_id)
678 max = MAX2(max, 2); /* VSPrimID */
679
680 /* GFX11: We prefer to compute RelAutoIndex using (WaveID * WaveSize + ThreadID).
681 * Older chips didn't have WaveID in LS.
682 * GFX12 doesn't have RelAutoIndex.
683 */
684 if (is_ls && sscreen->info.gfx_level <= GFX10_3)
685 max = MAX2(max, 1); /* RelAutoIndex */
686
687 return max;
688 }
689
si_shader_ls(struct si_screen * sscreen,struct si_shader * shader)690 static void si_shader_ls(struct si_screen *sscreen, struct si_shader *shader)
691 {
692 struct si_pm4_state *pm4;
693 uint64_t va;
694
695 assert(sscreen->info.gfx_level <= GFX8);
696
697 pm4 = si_get_shader_pm4_state(shader, NULL);
698 if (!pm4)
699 return;
700
701 va = shader->bo->gpu_address;
702 ac_pm4_set_reg(&pm4->base, R_00B520_SPI_SHADER_PGM_LO_LS, va >> 8);
703
704 shader->config.rsrc1 = S_00B528_VGPRS(si_shader_encode_vgprs(shader)) |
705 S_00B528_SGPRS(si_shader_encode_sgprs(shader)) |
706 S_00B528_VGPR_COMP_CNT(si_get_vs_vgpr_comp_cnt(sscreen, shader, false)) |
707 S_00B528_DX10_CLAMP(1) |
708 S_00B528_FLOAT_MODE(shader->config.float_mode);
709 shader->config.rsrc2 = S_00B52C_USER_SGPR(si_get_num_vs_user_sgprs(shader, SI_VS_NUM_USER_SGPR)) |
710 S_00B52C_SCRATCH_EN(shader->config.scratch_bytes_per_wave > 0);
711 ac_pm4_finalize(&pm4->base);
712 }
713
si_shader_hs(struct si_screen * sscreen,struct si_shader * shader)714 static void si_shader_hs(struct si_screen *sscreen, struct si_shader *shader)
715 {
716 struct si_pm4_state *pm4 = si_get_shader_pm4_state(shader, NULL);
717 if (!pm4)
718 return;
719
720 uint64_t va = shader->bo->gpu_address;
721 unsigned num_user_sgprs = sscreen->info.gfx_level >= GFX9 ?
722 si_get_num_vs_user_sgprs(shader, GFX9_TCS_NUM_USER_SGPR) :
723 GFX6_TCS_NUM_USER_SGPR;
724
725 if (sscreen->info.gfx_level >= GFX12) {
726 ac_pm4_set_reg(&pm4->base, R_00B420_SPI_SHADER_PGM_RSRC4_HS,
727 S_00B420_WAVE_LIMIT(0x3ff) |
728 S_00B420_GLG_FORCE_DISABLE(1) |
729 S_00B420_INST_PREF_SIZE(si_get_shader_prefetch_size(shader)));
730
731 ac_pm4_set_reg(&pm4->base, R_00B424_SPI_SHADER_PGM_LO_LS, va >> 8);
732 } else if (sscreen->info.gfx_level >= GFX11) {
733 ac_pm4_set_reg_idx3(&pm4->base, R_00B404_SPI_SHADER_PGM_RSRC4_HS,
734 ac_apply_cu_en(S_00B404_INST_PREF_SIZE(si_get_shader_prefetch_size(shader)) |
735 S_00B404_CU_EN(0xffff),
736 C_00B404_CU_EN, 16, &sscreen->info));
737
738 ac_pm4_set_reg(&pm4->base, R_00B520_SPI_SHADER_PGM_LO_LS, va >> 8);
739 } else if (sscreen->info.gfx_level >= GFX10) {
740 ac_pm4_set_reg(&pm4->base, R_00B520_SPI_SHADER_PGM_LO_LS, va >> 8);
741 } else if (sscreen->info.gfx_level >= GFX9) {
742 ac_pm4_set_reg(&pm4->base, R_00B410_SPI_SHADER_PGM_LO_LS, va >> 8);
743 } else {
744 ac_pm4_set_reg(&pm4->base, R_00B420_SPI_SHADER_PGM_LO_HS, va >> 8);
745 ac_pm4_set_reg(&pm4->base, R_00B424_SPI_SHADER_PGM_HI_HS,
746 S_00B424_MEM_BASE(sscreen->info.address32_hi >> 8));
747 }
748
749 ac_pm4_set_reg(&pm4->base, R_00B428_SPI_SHADER_PGM_RSRC1_HS,
750 S_00B428_VGPRS(si_shader_encode_vgprs(shader)) |
751 S_00B428_SGPRS(si_shader_encode_sgprs(shader)) |
752 S_00B428_DX10_CLAMP(sscreen->info.gfx_level < GFX12) |
753 S_00B428_MEM_ORDERED(si_shader_mem_ordered(shader)) |
754 S_00B428_FLOAT_MODE(shader->config.float_mode) |
755 S_00B428_LS_VGPR_COMP_CNT(sscreen->info.gfx_level >= GFX9 ?
756 si_get_vs_vgpr_comp_cnt(sscreen, shader, false) : 0));
757
758 shader->config.rsrc2 = S_00B42C_SCRATCH_EN(shader->config.scratch_bytes_per_wave > 0) |
759 S_00B42C_USER_SGPR(num_user_sgprs);
760
761 if (sscreen->info.gfx_level >= GFX10) {
762 shader->config.rsrc2 |= S_00B42C_USER_SGPR_MSB_GFX10(num_user_sgprs >> 5) |
763 S_00B42C_SHARED_VGPR_CNT(shader->config.num_shared_vgprs / 8);
764 } else if (sscreen->info.gfx_level >= GFX9) {
765 shader->config.rsrc2 |= S_00B42C_USER_SGPR_MSB_GFX9(num_user_sgprs >> 5);
766 } else {
767 shader->config.rsrc2 |= S_00B42C_OC_LDS_EN(1);
768 }
769
770 if (sscreen->info.gfx_level <= GFX8)
771 ac_pm4_set_reg(&pm4->base, R_00B42C_SPI_SHADER_PGM_RSRC2_HS, shader->config.rsrc2);
772
773 ac_pm4_finalize(&pm4->base);
774 }
775
si_emit_shader_es(struct si_context * sctx,unsigned index)776 static void si_emit_shader_es(struct si_context *sctx, unsigned index)
777 {
778 struct si_shader *shader = sctx->queued.named.es;
779
780 radeon_begin(&sctx->gfx_cs);
781 radeon_opt_set_context_reg(R_028AAC_VGT_ESGS_RING_ITEMSIZE,
782 SI_TRACKED_VGT_ESGS_RING_ITEMSIZE,
783 shader->selector->info.esgs_vertex_stride / 4);
784
785 if (shader->selector->stage == MESA_SHADER_TESS_EVAL)
786 radeon_opt_set_context_reg(R_028B6C_VGT_TF_PARAM, SI_TRACKED_VGT_TF_PARAM,
787 shader->vgt_tf_param);
788
789 if (shader->vgt_vertex_reuse_block_cntl)
790 radeon_opt_set_context_reg(R_028C58_VGT_VERTEX_REUSE_BLOCK_CNTL,
791 SI_TRACKED_VGT_VERTEX_REUSE_BLOCK_CNTL,
792 shader->vgt_vertex_reuse_block_cntl);
793 radeon_end_update_context_roll();
794 }
795
si_shader_es(struct si_screen * sscreen,struct si_shader * shader)796 static void si_shader_es(struct si_screen *sscreen, struct si_shader *shader)
797 {
798 struct si_pm4_state *pm4;
799 unsigned num_user_sgprs;
800 unsigned vgpr_comp_cnt;
801 uint64_t va;
802 unsigned oc_lds_en;
803
804 assert(sscreen->info.gfx_level <= GFX8);
805
806 pm4 = si_get_shader_pm4_state(shader, si_emit_shader_es);
807 if (!pm4)
808 return;
809
810 va = shader->bo->gpu_address;
811
812 if (shader->selector->stage == MESA_SHADER_VERTEX) {
813 vgpr_comp_cnt = si_get_vs_vgpr_comp_cnt(sscreen, shader, false);
814 num_user_sgprs = si_get_num_vs_user_sgprs(shader, SI_VS_NUM_USER_SGPR);
815 } else if (shader->selector->stage == MESA_SHADER_TESS_EVAL) {
816 vgpr_comp_cnt = shader->selector->info.uses_primid ? 3 : 2;
817 num_user_sgprs = SI_TES_NUM_USER_SGPR;
818 } else
819 unreachable("invalid shader selector type");
820
821 oc_lds_en = shader->selector->stage == MESA_SHADER_TESS_EVAL ? 1 : 0;
822
823 ac_pm4_set_reg(&pm4->base, R_00B320_SPI_SHADER_PGM_LO_ES, va >> 8);
824 ac_pm4_set_reg(&pm4->base, R_00B324_SPI_SHADER_PGM_HI_ES,
825 S_00B324_MEM_BASE(sscreen->info.address32_hi >> 8));
826 ac_pm4_set_reg(&pm4->base, R_00B328_SPI_SHADER_PGM_RSRC1_ES,
827 S_00B328_VGPRS(si_shader_encode_vgprs(shader)) |
828 S_00B328_SGPRS(si_shader_encode_sgprs(shader)) |
829 S_00B328_VGPR_COMP_CNT(vgpr_comp_cnt) |
830 S_00B328_DX10_CLAMP(1) |
831 S_00B328_FLOAT_MODE(shader->config.float_mode));
832 ac_pm4_set_reg(&pm4->base, R_00B32C_SPI_SHADER_PGM_RSRC2_ES,
833 S_00B32C_USER_SGPR(num_user_sgprs) | S_00B32C_OC_LDS_EN(oc_lds_en) |
834 S_00B32C_SCRATCH_EN(shader->config.scratch_bytes_per_wave > 0));
835
836 if (shader->selector->stage == MESA_SHADER_TESS_EVAL)
837 si_set_tesseval_regs(sscreen, shader->selector, shader);
838
839 polaris_set_vgt_vertex_reuse(sscreen, shader->selector, shader);
840 ac_pm4_finalize(&pm4->base);
841 }
842
gfx9_get_gs_info(struct si_shader_selector * es,struct si_shader_selector * gs,struct gfx9_gs_info * out)843 void gfx9_get_gs_info(struct si_shader_selector *es, struct si_shader_selector *gs,
844 struct gfx9_gs_info *out)
845 {
846 unsigned gs_num_invocations = MAX2(gs->info.base.gs.invocations, 1);
847 unsigned input_prim = gs->info.base.gs.input_primitive;
848 bool uses_adjacency = mesa_prim_has_adjacency((enum mesa_prim)input_prim);
849
850 /* All these are in dwords: */
851 /* We can't allow using the whole LDS, because GS waves compete with
852 * other shader stages for LDS space. */
853 const unsigned max_lds_size = 8 * 1024;
854 const unsigned esgs_itemsize = es->info.esgs_vertex_stride / 4;
855 unsigned esgs_lds_size;
856
857 /* All these are per subgroup: */
858 const unsigned max_out_prims = 32 * 1024;
859 const unsigned max_es_verts = 255;
860 const unsigned ideal_gs_prims = 64;
861 unsigned max_gs_prims, gs_prims;
862 unsigned min_es_verts, es_verts, worst_case_es_verts;
863
864 if (uses_adjacency || gs_num_invocations > 1)
865 max_gs_prims = 127 / gs_num_invocations;
866 else
867 max_gs_prims = 255;
868
869 /* MAX_PRIMS_PER_SUBGROUP = gs_prims * max_vert_out * gs_invocations.
870 * Make sure we don't go over the maximum value.
871 */
872 if (gs->info.base.gs.vertices_out > 0) {
873 max_gs_prims =
874 MIN2(max_gs_prims, max_out_prims / (gs->info.base.gs.vertices_out * gs_num_invocations));
875 }
876 assert(max_gs_prims > 0);
877
878 /* If the primitive has adjacency, halve the number of vertices
879 * that will be reused in multiple primitives.
880 */
881 min_es_verts = gs->info.gs_input_verts_per_prim / (uses_adjacency ? 2 : 1);
882
883 gs_prims = MIN2(ideal_gs_prims, max_gs_prims);
884 worst_case_es_verts = MIN2(min_es_verts * gs_prims, max_es_verts);
885
886 /* Compute ESGS LDS size based on the worst case number of ES vertices
887 * needed to create the target number of GS prims per subgroup.
888 */
889 esgs_lds_size = esgs_itemsize * worst_case_es_verts;
890
891 /* If total LDS usage is too big, refactor partitions based on ratio
892 * of ESGS item sizes.
893 */
894 if (esgs_lds_size > max_lds_size) {
895 /* Our target GS Prims Per Subgroup was too large. Calculate
896 * the maximum number of GS Prims Per Subgroup that will fit
897 * into LDS, capped by the maximum that the hardware can support.
898 */
899 gs_prims = MIN2((max_lds_size / (esgs_itemsize * min_es_verts)), max_gs_prims);
900 assert(gs_prims > 0);
901 worst_case_es_verts = MIN2(min_es_verts * gs_prims, max_es_verts);
902
903 esgs_lds_size = esgs_itemsize * worst_case_es_verts;
904 assert(esgs_lds_size <= max_lds_size);
905 }
906
907 /* Now calculate remaining ESGS information. */
908 if (esgs_lds_size)
909 es_verts = MIN2(esgs_lds_size / esgs_itemsize, max_es_verts);
910 else
911 es_verts = max_es_verts;
912
913 /* Vertices for adjacency primitives are not always reused, so restore
914 * it for ES_VERTS_PER_SUBGRP.
915 */
916 min_es_verts = gs->info.gs_input_verts_per_prim;
917
918 /* For normal primitives, the VGT only checks if they are past the ES
919 * verts per subgroup after allocating a full GS primitive and if they
920 * are, kick off a new subgroup. But if those additional ES verts are
921 * unique (e.g. not reused) we need to make sure there is enough LDS
922 * space to account for those ES verts beyond ES_VERTS_PER_SUBGRP.
923 */
924 es_verts -= min_es_verts - 1;
925
926 out->es_verts_per_subgroup = es_verts;
927 out->gs_prims_per_subgroup = gs_prims;
928 out->gs_inst_prims_in_subgroup = gs_prims * gs_num_invocations;
929 out->max_prims_per_subgroup = out->gs_inst_prims_in_subgroup * gs->info.base.gs.vertices_out;
930 out->esgs_ring_size = esgs_lds_size;
931
932 assert(out->max_prims_per_subgroup <= max_out_prims);
933 }
934
gfx9_set_gs_sgpr_num_es_outputs(struct si_context * sctx,unsigned esgs_vertex_stride)935 static void gfx9_set_gs_sgpr_num_es_outputs(struct si_context *sctx, unsigned esgs_vertex_stride)
936 {
937 /* The stride must always be odd (e.g. a multiple of 4 + 1) to reduce LDS bank conflicts. */
938 assert(!esgs_vertex_stride || esgs_vertex_stride % 4 == 1);
939 unsigned num_es_outputs = esgs_vertex_stride / 4;
940
941 /* If there are no ES outputs, GS doesn't use this SGPR field, so only set it if the number
942 * is non-zero.
943 */
944 if (num_es_outputs)
945 SET_FIELD(sctx->current_gs_state, GS_STATE_NUM_ES_OUTPUTS, num_es_outputs);
946 }
947
si_emit_shader_gs(struct si_context * sctx,unsigned index)948 static void si_emit_shader_gs(struct si_context *sctx, unsigned index)
949 {
950 struct si_shader *shader = sctx->queued.named.gs;
951
952 if (sctx->gfx_level >= GFX9)
953 gfx9_set_gs_sgpr_num_es_outputs(sctx, shader->key.ge.part.gs.es->info.esgs_vertex_stride / 4);
954
955 radeon_begin(&sctx->gfx_cs);
956
957 /* R_028A60_VGT_GSVS_RING_OFFSET_1, R_028A64_VGT_GSVS_RING_OFFSET_2
958 * R_028A68_VGT_GSVS_RING_OFFSET_3 */
959 radeon_opt_set_context_reg3(
960 R_028A60_VGT_GSVS_RING_OFFSET_1, SI_TRACKED_VGT_GSVS_RING_OFFSET_1,
961 shader->gs.vgt_gsvs_ring_offset_1, shader->gs.vgt_gsvs_ring_offset_2,
962 shader->gs.vgt_gsvs_ring_offset_3);
963
964 /* R_028AB0_VGT_GSVS_RING_ITEMSIZE */
965 radeon_opt_set_context_reg(R_028AB0_VGT_GSVS_RING_ITEMSIZE,
966 SI_TRACKED_VGT_GSVS_RING_ITEMSIZE,
967 shader->gs.vgt_gsvs_ring_itemsize);
968
969 /* R_028B38_VGT_GS_MAX_VERT_OUT */
970 radeon_opt_set_context_reg(R_028B38_VGT_GS_MAX_VERT_OUT, SI_TRACKED_VGT_GS_MAX_VERT_OUT,
971 shader->gs.vgt_gs_max_vert_out);
972
973 /* R_028B5C_VGT_GS_VERT_ITEMSIZE, R_028B60_VGT_GS_VERT_ITEMSIZE_1
974 * R_028B64_VGT_GS_VERT_ITEMSIZE_2, R_028B68_VGT_GS_VERT_ITEMSIZE_3 */
975 radeon_opt_set_context_reg4(
976 R_028B5C_VGT_GS_VERT_ITEMSIZE, SI_TRACKED_VGT_GS_VERT_ITEMSIZE,
977 shader->gs.vgt_gs_vert_itemsize, shader->gs.vgt_gs_vert_itemsize_1,
978 shader->gs.vgt_gs_vert_itemsize_2, shader->gs.vgt_gs_vert_itemsize_3);
979
980 /* R_028B90_VGT_GS_INSTANCE_CNT */
981 radeon_opt_set_context_reg(R_028B90_VGT_GS_INSTANCE_CNT, SI_TRACKED_VGT_GS_INSTANCE_CNT,
982 shader->gs.vgt_gs_instance_cnt);
983
984 if (sctx->gfx_level >= GFX9) {
985 /* R_028A44_VGT_GS_ONCHIP_CNTL */
986 radeon_opt_set_context_reg(R_028A44_VGT_GS_ONCHIP_CNTL, SI_TRACKED_VGT_GS_ONCHIP_CNTL,
987 shader->gs.vgt_gs_onchip_cntl);
988 /* R_028A94_VGT_GS_MAX_PRIMS_PER_SUBGROUP */
989 if (sctx->gfx_level == GFX9) {
990 radeon_opt_set_context_reg(R_028A94_VGT_GS_MAX_PRIMS_PER_SUBGROUP,
991 SI_TRACKED_VGT_GS_MAX_PRIMS_PER_SUBGROUP,
992 shader->gs.vgt_gs_max_prims_per_subgroup);
993 }
994
995 if (shader->key.ge.part.gs.es->stage == MESA_SHADER_TESS_EVAL)
996 radeon_opt_set_context_reg(R_028B6C_VGT_TF_PARAM, SI_TRACKED_VGT_TF_PARAM,
997 shader->vgt_tf_param);
998 if (shader->vgt_vertex_reuse_block_cntl)
999 radeon_opt_set_context_reg(R_028C58_VGT_VERTEX_REUSE_BLOCK_CNTL,
1000 SI_TRACKED_VGT_VERTEX_REUSE_BLOCK_CNTL,
1001 shader->vgt_vertex_reuse_block_cntl);
1002 }
1003 radeon_end_update_context_roll();
1004
1005 /* These don't cause any context rolls. */
1006 radeon_begin_again(&sctx->gfx_cs);
1007 if (sctx->gfx_level >= GFX7) {
1008 if (sctx->screen->info.uses_kernel_cu_mask) {
1009 radeon_opt_set_sh_reg_idx(R_00B21C_SPI_SHADER_PGM_RSRC3_GS,
1010 SI_TRACKED_SPI_SHADER_PGM_RSRC3_GS,
1011 3, shader->gs.spi_shader_pgm_rsrc3_gs);
1012 } else {
1013 radeon_opt_set_sh_reg(R_00B21C_SPI_SHADER_PGM_RSRC3_GS,
1014 SI_TRACKED_SPI_SHADER_PGM_RSRC3_GS,
1015 shader->gs.spi_shader_pgm_rsrc3_gs);
1016 }
1017 }
1018 if (sctx->gfx_level >= GFX10) {
1019 if (sctx->screen->info.uses_kernel_cu_mask) {
1020 radeon_opt_set_sh_reg_idx(R_00B204_SPI_SHADER_PGM_RSRC4_GS,
1021 SI_TRACKED_SPI_SHADER_PGM_RSRC4_GS,
1022 3, shader->gs.spi_shader_pgm_rsrc4_gs);
1023 } else {
1024 radeon_opt_set_sh_reg(R_00B204_SPI_SHADER_PGM_RSRC4_GS,
1025 SI_TRACKED_SPI_SHADER_PGM_RSRC4_GS,
1026 shader->gs.spi_shader_pgm_rsrc4_gs);
1027 }
1028 }
1029 radeon_end();
1030 }
1031
si_shader_gs(struct si_screen * sscreen,struct si_shader * shader)1032 static void si_shader_gs(struct si_screen *sscreen, struct si_shader *shader)
1033 {
1034 struct si_shader_selector *sel = shader->selector;
1035 const uint8_t *num_components = sel->info.num_stream_output_components;
1036 unsigned gs_num_invocations = sel->info.base.gs.invocations;
1037 struct si_pm4_state *pm4;
1038 uint64_t va;
1039 unsigned max_stream = util_last_bit(sel->info.base.gs.active_stream_mask);
1040 unsigned offset;
1041
1042 assert(sscreen->info.gfx_level < GFX11); /* gfx11 doesn't have the legacy pipeline */
1043
1044 pm4 = si_get_shader_pm4_state(shader, si_emit_shader_gs);
1045 if (!pm4)
1046 return;
1047
1048 offset = num_components[0] * sel->info.base.gs.vertices_out;
1049 shader->gs.vgt_gsvs_ring_offset_1 = offset;
1050
1051 if (max_stream >= 2)
1052 offset += num_components[1] * sel->info.base.gs.vertices_out;
1053 shader->gs.vgt_gsvs_ring_offset_2 = offset;
1054
1055 if (max_stream >= 3)
1056 offset += num_components[2] * sel->info.base.gs.vertices_out;
1057 shader->gs.vgt_gsvs_ring_offset_3 = offset;
1058
1059 if (max_stream >= 4)
1060 offset += num_components[3] * sel->info.base.gs.vertices_out;
1061 shader->gs.vgt_gsvs_ring_itemsize = offset;
1062
1063 /* The GSVS_RING_ITEMSIZE register takes 15 bits */
1064 assert(offset < (1 << 15));
1065
1066 shader->gs.vgt_gs_max_vert_out = sel->info.base.gs.vertices_out;
1067
1068 shader->gs.vgt_gs_vert_itemsize = num_components[0];
1069 shader->gs.vgt_gs_vert_itemsize_1 = (max_stream >= 2) ? num_components[1] : 0;
1070 shader->gs.vgt_gs_vert_itemsize_2 = (max_stream >= 3) ? num_components[2] : 0;
1071 shader->gs.vgt_gs_vert_itemsize_3 = (max_stream >= 4) ? num_components[3] : 0;
1072
1073 shader->gs.vgt_gs_instance_cnt =
1074 S_028B90_CNT(MIN2(gs_num_invocations, 127)) | S_028B90_ENABLE(gs_num_invocations > 0);
1075
1076 /* Copy over fields from the GS copy shader to make them easily accessible from GS. */
1077 shader->pa_cl_vs_out_cntl = shader->gs_copy_shader->pa_cl_vs_out_cntl;
1078
1079 va = shader->bo->gpu_address;
1080
1081 if (sscreen->info.gfx_level >= GFX9) {
1082 unsigned input_prim = sel->info.base.gs.input_primitive;
1083 gl_shader_stage es_stage = shader->key.ge.part.gs.es->stage;
1084 unsigned es_vgpr_comp_cnt, gs_vgpr_comp_cnt;
1085
1086 if (es_stage == MESA_SHADER_VERTEX) {
1087 es_vgpr_comp_cnt = si_get_vs_vgpr_comp_cnt(sscreen, shader, false);
1088 } else if (es_stage == MESA_SHADER_TESS_EVAL)
1089 es_vgpr_comp_cnt = shader->key.ge.part.gs.es->info.uses_primid ? 3 : 2;
1090 else
1091 unreachable("invalid shader selector type");
1092
1093 /* If offsets 4, 5 are used, GS_VGPR_COMP_CNT is ignored and
1094 * VGPR[0:4] are always loaded.
1095 */
1096 if (sel->info.uses_invocationid)
1097 gs_vgpr_comp_cnt = 3; /* VGPR3 contains InvocationID. */
1098 else if (sel->info.uses_primid)
1099 gs_vgpr_comp_cnt = 2; /* VGPR2 contains PrimitiveID. */
1100 else if (input_prim >= MESA_PRIM_TRIANGLES)
1101 gs_vgpr_comp_cnt = 1; /* VGPR1 contains offsets 2, 3 */
1102 else
1103 gs_vgpr_comp_cnt = 0; /* VGPR0 contains offsets 0, 1 */
1104
1105 unsigned num_user_sgprs;
1106 if (es_stage == MESA_SHADER_VERTEX)
1107 num_user_sgprs = si_get_num_vs_user_sgprs(shader, GFX9_GS_NUM_USER_SGPR);
1108 else
1109 num_user_sgprs = GFX9_GS_NUM_USER_SGPR;
1110
1111 if (sscreen->info.gfx_level >= GFX10) {
1112 ac_pm4_set_reg(&pm4->base, R_00B320_SPI_SHADER_PGM_LO_ES, va >> 8);
1113 } else {
1114 ac_pm4_set_reg(&pm4->base, R_00B210_SPI_SHADER_PGM_LO_ES, va >> 8);
1115 }
1116
1117 uint32_t rsrc1 = S_00B228_VGPRS(si_shader_encode_vgprs(shader)) |
1118 S_00B228_SGPRS(si_shader_encode_sgprs(shader)) |
1119 S_00B228_DX10_CLAMP(1) |
1120 S_00B228_MEM_ORDERED(si_shader_mem_ordered(shader)) |
1121 S_00B228_FLOAT_MODE(shader->config.float_mode) |
1122 S_00B228_GS_VGPR_COMP_CNT(gs_vgpr_comp_cnt);
1123 uint32_t rsrc2 = S_00B22C_USER_SGPR(num_user_sgprs) |
1124 S_00B22C_ES_VGPR_COMP_CNT(es_vgpr_comp_cnt) |
1125 S_00B22C_OC_LDS_EN(es_stage == MESA_SHADER_TESS_EVAL) |
1126 S_00B22C_LDS_SIZE(shader->config.lds_size) |
1127 S_00B22C_SCRATCH_EN(shader->config.scratch_bytes_per_wave > 0);
1128
1129 if (sscreen->info.gfx_level >= GFX10) {
1130 rsrc2 |= S_00B22C_USER_SGPR_MSB_GFX10(num_user_sgprs >> 5) |
1131 S_00B22C_SHARED_VGPR_CNT(shader->config.num_shared_vgprs / 8);
1132 } else {
1133 rsrc2 |= S_00B22C_USER_SGPR_MSB_GFX9(num_user_sgprs >> 5);
1134 }
1135
1136 ac_pm4_set_reg(&pm4->base, R_00B228_SPI_SHADER_PGM_RSRC1_GS, rsrc1);
1137 ac_pm4_set_reg(&pm4->base, R_00B22C_SPI_SHADER_PGM_RSRC2_GS, rsrc2);
1138
1139 shader->gs.spi_shader_pgm_rsrc3_gs =
1140 ac_apply_cu_en(S_00B21C_CU_EN(0xffff) |
1141 S_00B21C_WAVE_LIMIT(0x3F),
1142 C_00B21C_CU_EN, 0, &sscreen->info);
1143 shader->gs.spi_shader_pgm_rsrc4_gs =
1144 ac_apply_cu_en(S_00B204_CU_EN_GFX10(0xffff) |
1145 S_00B204_SPI_SHADER_LATE_ALLOC_GS_GFX10(0),
1146 C_00B204_CU_EN_GFX10, 16, &sscreen->info);
1147
1148 shader->gs.vgt_gs_onchip_cntl =
1149 S_028A44_ES_VERTS_PER_SUBGRP(shader->gs_info.es_verts_per_subgroup) |
1150 S_028A44_GS_PRIMS_PER_SUBGRP(shader->gs_info.gs_prims_per_subgroup) |
1151 S_028A44_GS_INST_PRIMS_IN_SUBGRP(shader->gs_info.gs_inst_prims_in_subgroup);
1152 shader->gs.vgt_gs_max_prims_per_subgroup =
1153 S_028A94_MAX_PRIMS_PER_SUBGROUP(shader->gs_info.max_prims_per_subgroup);
1154 shader->gs.vgt_esgs_ring_itemsize = shader->key.ge.part.gs.es->info.esgs_vertex_stride / 4;
1155
1156 if (es_stage == MESA_SHADER_TESS_EVAL)
1157 si_set_tesseval_regs(sscreen, shader->key.ge.part.gs.es, shader);
1158
1159 polaris_set_vgt_vertex_reuse(sscreen, shader->key.ge.part.gs.es, shader);
1160 } else {
1161 shader->gs.spi_shader_pgm_rsrc3_gs =
1162 ac_apply_cu_en(S_00B21C_CU_EN(0xffff) |
1163 S_00B21C_WAVE_LIMIT(0x3F),
1164 C_00B21C_CU_EN, 0, &sscreen->info);
1165
1166 ac_pm4_set_reg(&pm4->base, R_00B220_SPI_SHADER_PGM_LO_GS, va >> 8);
1167 ac_pm4_set_reg(&pm4->base, R_00B224_SPI_SHADER_PGM_HI_GS,
1168 S_00B224_MEM_BASE(sscreen->info.address32_hi >> 8));
1169
1170 ac_pm4_set_reg(&pm4->base, R_00B228_SPI_SHADER_PGM_RSRC1_GS,
1171 S_00B228_VGPRS(si_shader_encode_vgprs(shader)) |
1172 S_00B228_SGPRS(si_shader_encode_sgprs(shader)) |
1173 S_00B228_DX10_CLAMP(1) |
1174 S_00B228_FLOAT_MODE(shader->config.float_mode));
1175 ac_pm4_set_reg(&pm4->base, R_00B22C_SPI_SHADER_PGM_RSRC2_GS,
1176 S_00B22C_USER_SGPR(GFX6_GS_NUM_USER_SGPR) |
1177 S_00B22C_SCRATCH_EN(shader->config.scratch_bytes_per_wave > 0));
1178 }
1179 ac_pm4_finalize(&pm4->base);
1180 }
1181
gfx10_is_ngg_passthrough(struct si_shader * shader)1182 bool gfx10_is_ngg_passthrough(struct si_shader *shader)
1183 {
1184 struct si_shader_selector *sel = shader->selector;
1185
1186 /* Never use NGG passthrough if culling is possible even when it's not used by this shader,
1187 * so that we don't get context rolls when enabling and disabling NGG passthrough.
1188 */
1189 if (sel->screen->use_ngg_culling)
1190 return false;
1191
1192 /* The definition of NGG passthrough is:
1193 * - user GS is turned off (no amplification, no GS instancing, and no culling)
1194 * - VGT_ESGS_RING_ITEMSIZE is ignored (behaving as if it was equal to 1)
1195 * - vertex indices are packed into 1 VGPR
1196 * - Navi23 and later chips can optionally skip the gs_alloc_req message
1197 *
1198 * NGG passthrough still allows the use of LDS.
1199 */
1200 return sel->stage != MESA_SHADER_GEOMETRY && !si_shader_culling_enabled(shader);
1201 }
1202
1203 template <enum si_has_tess HAS_TESS>
gfx10_emit_shader_ngg(struct si_context * sctx,unsigned index)1204 static void gfx10_emit_shader_ngg(struct si_context *sctx, unsigned index)
1205 {
1206 struct si_shader *shader = sctx->queued.named.gs;
1207
1208 if (shader->selector->stage == MESA_SHADER_GEOMETRY)
1209 gfx9_set_gs_sgpr_num_es_outputs(sctx, shader->ngg.esgs_vertex_stride);
1210
1211 radeon_begin(&sctx->gfx_cs);
1212 if (HAS_TESS) {
1213 radeon_opt_set_context_reg(R_028B6C_VGT_TF_PARAM, SI_TRACKED_VGT_TF_PARAM,
1214 shader->vgt_tf_param);
1215 }
1216 radeon_opt_set_context_reg(R_0287FC_GE_MAX_OUTPUT_PER_SUBGROUP,
1217 SI_TRACKED_GE_MAX_OUTPUT_PER_SUBGROUP,
1218 shader->ngg.ge_max_output_per_subgroup);
1219 radeon_opt_set_context_reg(R_028B4C_GE_NGG_SUBGRP_CNTL, SI_TRACKED_GE_NGG_SUBGRP_CNTL,
1220 shader->ngg.ge_ngg_subgrp_cntl);
1221 radeon_opt_set_context_reg(R_028A84_VGT_PRIMITIVEID_EN, SI_TRACKED_VGT_PRIMITIVEID_EN,
1222 shader->ngg.vgt_primitiveid_en);
1223 if (sctx->gfx_level < GFX11) {
1224 radeon_opt_set_context_reg(R_028A44_VGT_GS_ONCHIP_CNTL, SI_TRACKED_VGT_GS_ONCHIP_CNTL,
1225 shader->ngg.vgt_gs_onchip_cntl);
1226 }
1227 radeon_opt_set_context_reg(R_028B38_VGT_GS_MAX_VERT_OUT, SI_TRACKED_VGT_GS_MAX_VERT_OUT,
1228 shader->ngg.vgt_gs_max_vert_out);
1229 radeon_opt_set_context_reg(R_028B90_VGT_GS_INSTANCE_CNT, SI_TRACKED_VGT_GS_INSTANCE_CNT,
1230 shader->ngg.vgt_gs_instance_cnt);
1231 radeon_opt_set_context_reg(R_0286C4_SPI_VS_OUT_CONFIG, SI_TRACKED_SPI_VS_OUT_CONFIG,
1232 shader->ngg.spi_vs_out_config);
1233 radeon_opt_set_context_reg(R_02870C_SPI_SHADER_POS_FORMAT,
1234 SI_TRACKED_SPI_SHADER_POS_FORMAT,
1235 shader->ngg.spi_shader_pos_format);
1236 radeon_opt_set_context_reg(R_028818_PA_CL_VTE_CNTL, SI_TRACKED_PA_CL_VTE_CNTL,
1237 shader->ngg.pa_cl_vte_cntl);
1238 radeon_end_update_context_roll();
1239
1240 /* These don't cause a context roll. */
1241 radeon_begin_again(&sctx->gfx_cs);
1242 if (sctx->screen->info.uses_kernel_cu_mask) {
1243 radeon_opt_set_sh_reg_idx(R_00B21C_SPI_SHADER_PGM_RSRC3_GS,
1244 SI_TRACKED_SPI_SHADER_PGM_RSRC3_GS,
1245 3, shader->ngg.spi_shader_pgm_rsrc3_gs);
1246 radeon_opt_set_sh_reg_idx(R_00B204_SPI_SHADER_PGM_RSRC4_GS,
1247 SI_TRACKED_SPI_SHADER_PGM_RSRC4_GS,
1248 3, shader->ngg.spi_shader_pgm_rsrc4_gs);
1249 } else {
1250 radeon_opt_set_sh_reg(R_00B21C_SPI_SHADER_PGM_RSRC3_GS,
1251 SI_TRACKED_SPI_SHADER_PGM_RSRC3_GS,
1252 shader->ngg.spi_shader_pgm_rsrc3_gs);
1253 radeon_opt_set_sh_reg(R_00B204_SPI_SHADER_PGM_RSRC4_GS,
1254 SI_TRACKED_SPI_SHADER_PGM_RSRC4_GS,
1255 shader->ngg.spi_shader_pgm_rsrc4_gs);
1256 }
1257 radeon_opt_set_uconfig_reg(R_030980_GE_PC_ALLOC, SI_TRACKED_GE_PC_ALLOC,
1258 shader->ngg.ge_pc_alloc);
1259 radeon_end();
1260 }
1261
1262 template <enum si_has_tess HAS_TESS>
gfx11_dgpu_emit_shader_ngg(struct si_context * sctx,unsigned index)1263 static void gfx11_dgpu_emit_shader_ngg(struct si_context *sctx, unsigned index)
1264 {
1265 struct si_shader *shader = sctx->queued.named.gs;
1266
1267 if (shader->selector->stage == MESA_SHADER_GEOMETRY)
1268 gfx9_set_gs_sgpr_num_es_outputs(sctx, shader->ngg.esgs_vertex_stride);
1269
1270 radeon_begin(&sctx->gfx_cs);
1271 gfx11_begin_packed_context_regs();
1272 if (HAS_TESS) {
1273 gfx11_opt_set_context_reg(R_028B6C_VGT_TF_PARAM, SI_TRACKED_VGT_TF_PARAM,
1274 shader->vgt_tf_param);
1275 }
1276 gfx11_opt_set_context_reg(R_0287FC_GE_MAX_OUTPUT_PER_SUBGROUP,
1277 SI_TRACKED_GE_MAX_OUTPUT_PER_SUBGROUP,
1278 shader->ngg.ge_max_output_per_subgroup);
1279 gfx11_opt_set_context_reg(R_028B4C_GE_NGG_SUBGRP_CNTL, SI_TRACKED_GE_NGG_SUBGRP_CNTL,
1280 shader->ngg.ge_ngg_subgrp_cntl);
1281 gfx11_opt_set_context_reg(R_028A84_VGT_PRIMITIVEID_EN, SI_TRACKED_VGT_PRIMITIVEID_EN,
1282 shader->ngg.vgt_primitiveid_en);
1283 gfx11_opt_set_context_reg(R_028B38_VGT_GS_MAX_VERT_OUT, SI_TRACKED_VGT_GS_MAX_VERT_OUT,
1284 shader->ngg.vgt_gs_max_vert_out);
1285 gfx11_opt_set_context_reg(R_028B90_VGT_GS_INSTANCE_CNT, SI_TRACKED_VGT_GS_INSTANCE_CNT,
1286 shader->ngg.vgt_gs_instance_cnt);
1287 gfx11_opt_set_context_reg(R_0286C4_SPI_VS_OUT_CONFIG, SI_TRACKED_SPI_VS_OUT_CONFIG,
1288 shader->ngg.spi_vs_out_config);
1289 gfx11_opt_set_context_reg(R_02870C_SPI_SHADER_POS_FORMAT, SI_TRACKED_SPI_SHADER_POS_FORMAT,
1290 shader->ngg.spi_shader_pos_format);
1291 gfx11_opt_set_context_reg(R_028818_PA_CL_VTE_CNTL, SI_TRACKED_PA_CL_VTE_CNTL,
1292 shader->ngg.pa_cl_vte_cntl);
1293 gfx11_end_packed_context_regs();
1294
1295 assert(!sctx->screen->info.uses_kernel_cu_mask);
1296 if (sctx->screen->info.has_set_sh_pairs_packed) {
1297 gfx11_opt_push_gfx_sh_reg(R_00B21C_SPI_SHADER_PGM_RSRC3_GS,
1298 SI_TRACKED_SPI_SHADER_PGM_RSRC3_GS,
1299 shader->gs.spi_shader_pgm_rsrc3_gs);
1300 gfx11_opt_push_gfx_sh_reg(R_00B204_SPI_SHADER_PGM_RSRC4_GS,
1301 SI_TRACKED_SPI_SHADER_PGM_RSRC4_GS,
1302 shader->gs.spi_shader_pgm_rsrc4_gs);
1303 } else {
1304 if (sctx->screen->info.uses_kernel_cu_mask) {
1305 radeon_opt_set_sh_reg_idx(R_00B21C_SPI_SHADER_PGM_RSRC3_GS,
1306 SI_TRACKED_SPI_SHADER_PGM_RSRC3_GS,
1307 3, shader->ngg.spi_shader_pgm_rsrc3_gs);
1308 radeon_opt_set_sh_reg_idx(R_00B204_SPI_SHADER_PGM_RSRC4_GS,
1309 SI_TRACKED_SPI_SHADER_PGM_RSRC4_GS,
1310 3, shader->ngg.spi_shader_pgm_rsrc4_gs);
1311 } else {
1312 radeon_opt_set_sh_reg(R_00B21C_SPI_SHADER_PGM_RSRC3_GS,
1313 SI_TRACKED_SPI_SHADER_PGM_RSRC3_GS,
1314 shader->ngg.spi_shader_pgm_rsrc3_gs);
1315 radeon_opt_set_sh_reg(R_00B204_SPI_SHADER_PGM_RSRC4_GS,
1316 SI_TRACKED_SPI_SHADER_PGM_RSRC4_GS,
1317 shader->ngg.spi_shader_pgm_rsrc4_gs);
1318 }
1319 }
1320
1321 radeon_opt_set_uconfig_reg(R_030980_GE_PC_ALLOC, SI_TRACKED_GE_PC_ALLOC,
1322 shader->ngg.ge_pc_alloc);
1323 radeon_end();
1324 }
1325
1326 template <enum si_has_tess HAS_TESS>
gfx12_emit_shader_ngg(struct si_context * sctx,unsigned index)1327 static void gfx12_emit_shader_ngg(struct si_context *sctx, unsigned index)
1328 {
1329 struct si_shader *shader = sctx->queued.named.gs;
1330
1331 if (shader->selector->stage == MESA_SHADER_GEOMETRY)
1332 gfx9_set_gs_sgpr_num_es_outputs(sctx, shader->ngg.esgs_vertex_stride);
1333
1334 radeon_begin(&sctx->gfx_cs);
1335 gfx12_begin_context_regs();
1336 if (HAS_TESS) {
1337 gfx12_opt_set_context_reg(R_028AA4_VGT_TF_PARAM, SI_TRACKED_VGT_TF_PARAM,
1338 shader->vgt_tf_param);
1339 }
1340 gfx12_opt_set_context_reg(R_0287FC_GE_MAX_OUTPUT_PER_SUBGROUP,
1341 SI_TRACKED_GE_MAX_OUTPUT_PER_SUBGROUP,
1342 shader->ngg.ge_max_output_per_subgroup);
1343 gfx12_opt_set_context_reg(R_028B4C_GE_NGG_SUBGRP_CNTL, SI_TRACKED_GE_NGG_SUBGRP_CNTL,
1344 shader->ngg.ge_ngg_subgrp_cntl);
1345 gfx12_opt_set_context_reg(R_028B38_VGT_GS_MAX_VERT_OUT, SI_TRACKED_VGT_GS_MAX_VERT_OUT,
1346 shader->ngg.vgt_gs_max_vert_out);
1347 gfx12_opt_set_context_reg(R_028B3C_VGT_GS_INSTANCE_CNT, SI_TRACKED_VGT_GS_INSTANCE_CNT,
1348 shader->ngg.vgt_gs_instance_cnt);
1349 gfx12_opt_set_context_reg(R_02864C_SPI_SHADER_POS_FORMAT, SI_TRACKED_SPI_SHADER_POS_FORMAT,
1350 shader->ngg.spi_shader_pos_format);
1351 gfx12_opt_set_context_reg(R_028814_PA_CL_VTE_CNTL, SI_TRACKED_PA_CL_VTE_CNTL,
1352 shader->ngg.pa_cl_vte_cntl);
1353 gfx12_end_context_regs();
1354
1355 radeon_opt_set_uconfig_reg(R_030988_VGT_PRIMITIVEID_EN,
1356 SI_TRACKED_VGT_PRIMITIVEID_EN_UCONFIG,
1357 shader->ngg.vgt_primitiveid_en);
1358 radeon_end(); /* don't track context rolls on GFX12 */
1359
1360 assert(!sctx->screen->info.uses_kernel_cu_mask);
1361 gfx12_opt_push_gfx_sh_reg(R_00B220_SPI_SHADER_PGM_RSRC4_GS,
1362 SI_TRACKED_SPI_SHADER_PGM_RSRC4_GS,
1363 shader->ngg.spi_shader_pgm_rsrc4_gs);
1364 }
1365
si_get_input_prim(const struct si_shader_selector * gs,const union si_shader_key * key,bool return_unknown)1366 unsigned si_get_input_prim(const struct si_shader_selector *gs, const union si_shader_key *key,
1367 bool return_unknown)
1368 {
1369 if (gs->stage == MESA_SHADER_GEOMETRY)
1370 return gs->info.base.gs.input_primitive;
1371
1372 if (gs->stage == MESA_SHADER_TESS_EVAL) {
1373 if (gs->info.base.tess.point_mode)
1374 return MESA_PRIM_POINTS;
1375 if (gs->info.base.tess._primitive_mode == TESS_PRIMITIVE_ISOLINES)
1376 return MESA_PRIM_LINES;
1377 return MESA_PRIM_TRIANGLES;
1378 }
1379
1380 assert(gs->stage == MESA_SHADER_VERTEX);
1381
1382 if (key->ge.opt.ngg_culling & SI_NGG_CULL_VS_LINES)
1383 return MESA_PRIM_LINES;
1384
1385 switch (key->ge.opt.ngg_vs_streamout_num_verts_per_prim) {
1386 case 3:
1387 return MESA_PRIM_TRIANGLES;
1388 case 2:
1389 return MESA_PRIM_LINES;
1390 case 1:
1391 return MESA_PRIM_POINTS;
1392 }
1393
1394 if (return_unknown)
1395 return MESA_PRIM_UNKNOWN;
1396 else
1397 return MESA_PRIM_TRIANGLES; /* worst case for all callers */
1398 }
1399
1400 /* Return a simplified primitive type, e.g. don't return *_STRIP and *_FAN.
1401 * This returns MESA_PRIM_UNKNOWN if the primitive type is not known at compile time.
1402 */
si_get_output_prim_simplified(const struct si_shader_selector * sel,const union si_shader_key * key)1403 unsigned si_get_output_prim_simplified(const struct si_shader_selector *sel,
1404 const union si_shader_key *key)
1405 {
1406 if (sel->stage == MESA_SHADER_GEOMETRY) {
1407 if (util_rast_prim_is_triangles(sel->info.base.gs.output_primitive))
1408 return MESA_PRIM_TRIANGLES;
1409 else if (util_prim_is_lines(sel->info.base.gs.output_primitive))
1410 return MESA_PRIM_LINES;
1411 else
1412 return MESA_PRIM_POINTS;
1413 }
1414
1415 if (sel->stage == MESA_SHADER_VERTEX && sel->info.base.vs.blit_sgprs_amd)
1416 return SI_PRIM_RECTANGLE_LIST;
1417
1418 /* It's the same as the input primitive type for VS and TES. */
1419 return si_get_input_prim(sel, key, true);
1420 }
1421
si_get_num_vertices_per_output_prim(struct si_shader * shader)1422 unsigned si_get_num_vertices_per_output_prim(struct si_shader *shader)
1423 {
1424 unsigned prim = si_get_output_prim_simplified(shader->selector, &shader->key);
1425
1426 switch (prim) {
1427 case MESA_PRIM_TRIANGLES:
1428 case SI_PRIM_RECTANGLE_LIST:
1429 return 3;
1430 case MESA_PRIM_LINES:
1431 return 2;
1432 case MESA_PRIM_POINTS:
1433 return 1;
1434 case MESA_PRIM_UNKNOWN:
1435 return 0;
1436 default:
1437 unreachable("unexpected prim type");
1438 }
1439 }
1440
si_get_vs_out_cntl(const struct si_shader_selector * sel,const struct si_shader * shader,bool ngg)1441 static unsigned si_get_vs_out_cntl(const struct si_shader_selector *sel,
1442 const struct si_shader *shader, bool ngg)
1443 {
1444 /* Clip distances can be killed, but cull distances can't. */
1445 unsigned clipcull_mask = (sel->info.clipdist_mask & ~shader->key.ge.opt.kill_clip_distances) |
1446 sel->info.culldist_mask;
1447 bool writes_psize = sel->info.writes_psize && !shader->key.ge.opt.kill_pointsize;
1448 bool writes_layer = sel->info.writes_layer && !shader->key.ge.opt.kill_layer;
1449 bool misc_vec_ena = writes_psize || (sel->info.writes_edgeflag && !ngg) ||
1450 writes_layer || sel->info.writes_viewport_index ||
1451 sel->screen->options.vrs2x2;
1452
1453 return S_02881C_VS_OUT_CCDIST0_VEC_ENA((clipcull_mask & 0x0F) != 0) |
1454 S_02881C_VS_OUT_CCDIST1_VEC_ENA((clipcull_mask & 0xF0) != 0) |
1455 S_02881C_USE_VTX_POINT_SIZE(writes_psize) |
1456 S_02881C_USE_VTX_EDGE_FLAG(sel->info.writes_edgeflag && !ngg) |
1457 S_02881C_USE_VTX_VRS_RATE(sel->screen->options.vrs2x2) |
1458 S_02881C_USE_VTX_RENDER_TARGET_INDX(writes_layer) |
1459 S_02881C_USE_VTX_VIEWPORT_INDX(sel->info.writes_viewport_index) |
1460 S_02881C_VS_OUT_MISC_VEC_ENA(misc_vec_ena) |
1461 S_02881C_VS_OUT_MISC_SIDE_BUS_ENA(misc_vec_ena ||
1462 (sel->screen->info.gfx_level >= GFX10_3 &&
1463 shader->info.nr_pos_exports > 1));
1464 }
1465
1466 /* Return the number of allocated param exports. This can be more than the number of param
1467 * exports in the shader.
1468 */
si_shader_num_alloc_param_exports(struct si_shader * shader)1469 unsigned si_shader_num_alloc_param_exports(struct si_shader *shader)
1470 {
1471 unsigned num_params = shader->info.nr_param_exports;
1472
1473 /* Since there is no alloc/dealloc mechanism for the 12-bit ordered IDs on GFX12, they can wrap
1474 * around if there are more than 2^12 workgroups, causing 2 workgroups to get the same
1475 * ordered ID, which can deadlock the "ordered add" loop.
1476 *
1477 * The recommended solution is to use the alloc/dealloc mechanism of the attribute ring to limit
1478 * the number of workgroups in flight and thus the number of ordered IDs in flight.
1479 */
1480 if (shader->selector->screen->info.gfx_level >= GFX12 && si_shader_uses_streamout(shader))
1481 num_params = MAX2(num_params, 8);
1482
1483 return num_params;
1484 }
1485
1486 /**
1487 * Prepare the PM4 image for \p shader, which will run as a merged ESGS shader
1488 * in NGG mode.
1489 */
gfx10_shader_ngg(struct si_screen * sscreen,struct si_shader * shader)1490 static void gfx10_shader_ngg(struct si_screen *sscreen, struct si_shader *shader)
1491 {
1492 const struct si_shader_selector *gs_sel = shader->selector;
1493 const struct si_shader_info *gs_info = &gs_sel->info;
1494 const gl_shader_stage gs_stage = shader->selector->stage;
1495 const struct si_shader_selector *es_sel =
1496 shader->previous_stage_sel ? shader->previous_stage_sel : shader->selector;
1497 const struct si_shader_info *es_info = &es_sel->info;
1498 const gl_shader_stage es_stage = es_sel->stage;
1499 unsigned num_user_sgprs;
1500 unsigned es_vgpr_comp_cnt, gs_vgpr_comp_cnt;
1501 uint64_t va;
1502 bool window_space = gs_sel->stage == MESA_SHADER_VERTEX ?
1503 gs_info->base.vs.window_space_position : 0;
1504 bool es_enable_prim_id = shader->key.ge.mono.u.vs_export_prim_id || es_info->uses_primid;
1505 unsigned gs_num_invocations = gs_sel->stage == MESA_SHADER_GEOMETRY ?
1506 CLAMP(gs_info->base.gs.invocations, 1, 32) : 0;
1507 unsigned input_prim = si_get_input_prim(gs_sel, &shader->key, false);
1508
1509 struct si_pm4_state *pm4 = si_get_shader_pm4_state(shader, NULL);
1510 if (!pm4)
1511 return;
1512
1513 if (sscreen->info.gfx_level >= GFX12) {
1514 if (es_stage == MESA_SHADER_TESS_EVAL)
1515 pm4->atom.emit = gfx12_emit_shader_ngg<TESS_ON>;
1516 else
1517 pm4->atom.emit = gfx12_emit_shader_ngg<TESS_OFF>;
1518 } else if (sscreen->info.has_set_context_pairs_packed) {
1519 if (es_stage == MESA_SHADER_TESS_EVAL)
1520 pm4->atom.emit = gfx11_dgpu_emit_shader_ngg<TESS_ON>;
1521 else
1522 pm4->atom.emit = gfx11_dgpu_emit_shader_ngg<TESS_OFF>;
1523 } else {
1524 if (es_stage == MESA_SHADER_TESS_EVAL)
1525 pm4->atom.emit = gfx10_emit_shader_ngg<TESS_ON>;
1526 else
1527 pm4->atom.emit = gfx10_emit_shader_ngg<TESS_OFF>;
1528 }
1529
1530 va = shader->bo->gpu_address;
1531
1532 if (es_stage == MESA_SHADER_VERTEX) {
1533 es_vgpr_comp_cnt = si_get_vs_vgpr_comp_cnt(sscreen, shader, false);
1534
1535 if (es_info->base.vs.blit_sgprs_amd) {
1536 num_user_sgprs =
1537 SI_SGPR_VS_BLIT_DATA + es_info->base.vs.blit_sgprs_amd;
1538 } else {
1539 num_user_sgprs = si_get_num_vs_user_sgprs(shader, GFX9_GS_NUM_USER_SGPR);
1540 }
1541 } else {
1542 assert(es_stage == MESA_SHADER_TESS_EVAL);
1543 es_vgpr_comp_cnt = es_enable_prim_id ? 3 : 2;
1544 num_user_sgprs = GFX9_GS_NUM_USER_SGPR;
1545 }
1546
1547 /* Primitives with adjancency can only occur without tessellation. */
1548 assert(gs_info->gs_input_verts_per_prim <= 3 || es_stage == MESA_SHADER_VERTEX);
1549
1550 if (sscreen->info.gfx_level >= GFX12) {
1551 if (gs_info->gs_input_verts_per_prim >= 4)
1552 gs_vgpr_comp_cnt = 2; /* VGPR2 contains offsets 3-5 */
1553 else if ((gs_stage == MESA_SHADER_GEOMETRY && gs_info->uses_primid) ||
1554 (gs_stage == MESA_SHADER_VERTEX && shader->key.ge.mono.u.vs_export_prim_id))
1555 gs_vgpr_comp_cnt = 1; /* VGPR1 contains PrimitiveID */
1556 else
1557 gs_vgpr_comp_cnt = 0; /* VGPR0 contains offsets 0-2, edgeflags, GS invocation ID. */
1558 } else {
1559 /* If offsets 4, 5 are used, GS_VGPR_COMP_CNT is ignored and
1560 * VGPR[0:4] are always loaded.
1561 *
1562 * Vertex shaders always need to load VGPR3, because they need to
1563 * pass edge flags for decomposed primitives (such as quads) to the PA
1564 * for the GL_LINE polygon mode to skip rendering lines on inner edges.
1565 */
1566 if (gs_info->uses_invocationid ||
1567 (gfx10_has_variable_edgeflags(shader) && !gfx10_is_ngg_passthrough(shader)))
1568 gs_vgpr_comp_cnt = 3; /* VGPR3 contains InvocationID, edge flags. */
1569 else if ((gs_stage == MESA_SHADER_GEOMETRY && gs_info->uses_primid) ||
1570 (gs_stage == MESA_SHADER_VERTEX && shader->key.ge.mono.u.vs_export_prim_id))
1571 gs_vgpr_comp_cnt = 2; /* VGPR2 contains PrimitiveID. */
1572 else if (input_prim >= MESA_PRIM_TRIANGLES && !gfx10_is_ngg_passthrough(shader))
1573 gs_vgpr_comp_cnt = 1; /* VGPR1 contains offsets 2, 3 */
1574 else
1575 gs_vgpr_comp_cnt = 0; /* VGPR0 contains offsets 0, 1 */
1576 }
1577
1578 if (sscreen->info.gfx_level >= GFX12) {
1579 ac_pm4_set_reg(&pm4->base, R_00B224_SPI_SHADER_PGM_LO_ES, va >> 8);
1580 } else {
1581 ac_pm4_set_reg(&pm4->base, R_00B320_SPI_SHADER_PGM_LO_ES, va >> 8);
1582 }
1583
1584 ac_pm4_set_reg(&pm4->base, R_00B228_SPI_SHADER_PGM_RSRC1_GS,
1585 S_00B228_VGPRS(si_shader_encode_vgprs(shader)) |
1586 S_00B228_FLOAT_MODE(shader->config.float_mode) |
1587 S_00B228_DX10_CLAMP(sscreen->info.gfx_level < GFX12) |
1588 S_00B228_MEM_ORDERED(si_shader_mem_ordered(shader)) |
1589 S_00B228_GS_VGPR_COMP_CNT(gs_vgpr_comp_cnt));
1590 ac_pm4_set_reg(&pm4->base, R_00B22C_SPI_SHADER_PGM_RSRC2_GS,
1591 S_00B22C_SCRATCH_EN(shader->config.scratch_bytes_per_wave > 0) |
1592 S_00B22C_USER_SGPR(num_user_sgprs) |
1593 S_00B22C_ES_VGPR_COMP_CNT(es_vgpr_comp_cnt) |
1594 S_00B22C_OC_LDS_EN(es_stage == MESA_SHADER_TESS_EVAL) |
1595 S_00B22C_LDS_SIZE(shader->config.lds_size) |
1596 S_00B22C_USER_SGPR_MSB_GFX10(num_user_sgprs >> 5) |
1597 S_00B22C_SHARED_VGPR_CNT(shader->config.num_shared_vgprs / 8));
1598
1599 /* Set register values emitted conditionally in gfx10_emit_shader_ngg_*. */
1600 shader->ngg.spi_shader_pos_format =
1601 S_02870C_POS0_EXPORT_FORMAT(V_02870C_SPI_SHADER_4COMP) |
1602 S_02870C_POS1_EXPORT_FORMAT(shader->info.nr_pos_exports > 1 ? V_02870C_SPI_SHADER_4COMP
1603 : V_02870C_SPI_SHADER_NONE) |
1604 S_02870C_POS2_EXPORT_FORMAT(shader->info.nr_pos_exports > 2 ? V_02870C_SPI_SHADER_4COMP
1605 : V_02870C_SPI_SHADER_NONE) |
1606 S_02870C_POS3_EXPORT_FORMAT(shader->info.nr_pos_exports > 3 ? V_02870C_SPI_SHADER_4COMP
1607 : V_02870C_SPI_SHADER_NONE);
1608 shader->ngg.ge_max_output_per_subgroup = S_0287FC_MAX_VERTS_PER_SUBGROUP(shader->ngg.max_out_verts);
1609 shader->ngg.vgt_gs_instance_cnt =
1610 S_028B90_ENABLE(gs_num_invocations > 1) |
1611 S_028B90_CNT(gs_num_invocations) |
1612 S_028B90_EN_MAX_VERT_OUT_PER_GS_INSTANCE(shader->ngg.max_vert_out_per_gs_instance);
1613 shader->pa_cl_vs_out_cntl = si_get_vs_out_cntl(shader->selector, shader, true);
1614
1615 if (gs_stage == MESA_SHADER_GEOMETRY) {
1616 shader->ngg.esgs_vertex_stride = es_sel->info.esgs_vertex_stride / 4;
1617 shader->ngg.vgt_gs_max_vert_out = gs_sel->info.base.gs.vertices_out;
1618 shader->ngg.ge_ngg_subgrp_cntl = S_028B4C_PRIM_AMP_FACTOR(gs_sel->info.base.gs.vertices_out);
1619 } else {
1620 shader->ngg.esgs_vertex_stride = 1;
1621 shader->ngg.vgt_gs_max_vert_out = 1;
1622 shader->ngg.ge_ngg_subgrp_cntl = S_028B4C_PRIM_AMP_FACTOR(1);
1623 }
1624
1625 if (es_stage == MESA_SHADER_TESS_EVAL)
1626 si_set_tesseval_regs(sscreen, es_sel, shader);
1627
1628 shader->ngg.vgt_primitiveid_en =
1629 S_028A84_NGG_DISABLE_PROVOK_REUSE(shader->key.ge.mono.u.vs_export_prim_id ||
1630 gs_sel->info.writes_primid);
1631
1632 if (sscreen->info.gfx_level >= GFX12) {
1633 unsigned num_params = si_shader_num_alloc_param_exports(shader);
1634 unsigned wave_limit_per_se = 0x3ff;
1635
1636 /* This tuning adds up to 50% streamout performance. */
1637 if (si_shader_uses_streamout(shader)) {
1638 unsigned num_streamout_vec4s = DIV_ROUND_UP(shader->selector->info.num_streamout_components, 4);
1639
1640 /* TODO: Tested on a pre-production chip. Re-test on the final chip. */
1641 if (num_streamout_vec4s <= 4)
1642 wave_limit_per_se = 48;
1643 else if (num_streamout_vec4s <= 5)
1644 wave_limit_per_se = 24;
1645 else if (num_streamout_vec4s <= 6)
1646 wave_limit_per_se = 20;
1647 else if (num_streamout_vec4s <= 8)
1648 wave_limit_per_se = 18;
1649 else if (num_streamout_vec4s <= 11)
1650 wave_limit_per_se = 17;
1651 else if (num_streamout_vec4s <= 12)
1652 wave_limit_per_se = 16;
1653 else if (num_streamout_vec4s <= 15)
1654 wave_limit_per_se = 15;
1655 else
1656 wave_limit_per_se = 14;
1657 }
1658
1659 shader->ngg.spi_shader_pgm_rsrc4_gs = S_00B220_SPI_SHADER_LATE_ALLOC_GS(127) |
1660 S_00B220_GLG_FORCE_DISABLE(1) |
1661 S_00B220_WAVE_LIMIT(wave_limit_per_se) |
1662 S_00B220_INST_PREF_SIZE(si_get_shader_prefetch_size(shader));
1663 shader->ngg.spi_vs_out_config = S_00B0C4_VS_EXPORT_COUNT(MAX2(num_params, 1) - 1) |
1664 S_00B0C4_NO_PC_EXPORT(num_params == 0);
1665 } else {
1666 unsigned late_alloc_wave64, cu_mask;
1667
1668 ac_compute_late_alloc(&sscreen->info, true, si_shader_culling_enabled(shader),
1669 shader->config.scratch_bytes_per_wave > 0,
1670 &late_alloc_wave64, &cu_mask);
1671
1672 /* Oversubscribe PC. This improves performance when there are too many varyings. */
1673 unsigned oversub_pc_lines, oversub_pc_factor = 1;
1674
1675 if (si_shader_culling_enabled(shader)) {
1676 /* Be more aggressive with NGG culling. */
1677 if (shader->info.nr_param_exports > 4)
1678 oversub_pc_factor = 4;
1679 else if (shader->info.nr_param_exports > 2)
1680 oversub_pc_factor = 3;
1681 else
1682 oversub_pc_factor = 2;
1683 }
1684 oversub_pc_lines = late_alloc_wave64 ? (sscreen->info.pc_lines / 4) * oversub_pc_factor : 0;
1685 shader->ngg.ge_pc_alloc = S_030980_OVERSUB_EN(oversub_pc_lines > 0) |
1686 S_030980_NUM_PC_LINES(oversub_pc_lines - 1);
1687 shader->ngg.vgt_primitiveid_en |= S_028A84_PRIMITIVEID_EN(es_enable_prim_id);
1688 shader->ngg.spi_shader_pgm_rsrc3_gs =
1689 ac_apply_cu_en(S_00B21C_CU_EN(cu_mask) |
1690 S_00B21C_WAVE_LIMIT(0x3F),
1691 C_00B21C_CU_EN, 0, &sscreen->info);
1692 shader->ngg.spi_shader_pgm_rsrc4_gs = S_00B204_SPI_SHADER_LATE_ALLOC_GS_GFX10(late_alloc_wave64);
1693 shader->ngg.spi_vs_out_config =
1694 S_0286C4_VS_EXPORT_COUNT(MAX2(shader->info.nr_param_exports, 1) - 1) |
1695 S_0286C4_NO_PC_EXPORT(shader->info.nr_param_exports == 0);
1696
1697 if (sscreen->info.gfx_level >= GFX11) {
1698 shader->ngg.spi_shader_pgm_rsrc4_gs |=
1699 ac_apply_cu_en(S_00B204_CU_EN_GFX11(0x1) |
1700 S_00B204_INST_PREF_SIZE(si_get_shader_prefetch_size(shader)),
1701 C_00B204_CU_EN_GFX11, 16, &sscreen->info);
1702 } else {
1703 shader->ngg.spi_shader_pgm_rsrc4_gs |=
1704 ac_apply_cu_en(S_00B204_CU_EN_GFX10(0xffff),
1705 C_00B204_CU_EN_GFX10, 16, &sscreen->info);
1706 }
1707 }
1708
1709 if (sscreen->info.gfx_level >= GFX11) {
1710 /* This should be <= 252 for performance on Gfx11. 256 works too but is slower. */
1711 unsigned max_prim_grp_size = sscreen->info.gfx_level >= GFX12 ? 256 : 252;
1712 unsigned prim_amp_factor = gs_stage == MESA_SHADER_GEOMETRY ?
1713 gs_sel->info.base.gs.vertices_out : 1;
1714
1715 shader->ge_cntl = S_03096C_PRIMS_PER_SUBGRP(shader->ngg.max_gsprims) |
1716 S_03096C_VERTS_PER_SUBGRP(shader->ngg.hw_max_esverts) |
1717 S_03096C_PRIM_GRP_SIZE_GFX11(
1718 CLAMP(max_prim_grp_size / MAX2(prim_amp_factor, 1), 1, 256)) |
1719 S_03096C_DIS_PG_SIZE_ADJUST_FOR_STRIP(sscreen->info.gfx_level >= GFX12);
1720 } else {
1721 shader->ge_cntl = S_03096C_PRIM_GRP_SIZE_GFX10(shader->ngg.max_gsprims) |
1722 S_03096C_VERT_GRP_SIZE(shader->ngg.hw_max_esverts);
1723
1724 shader->ngg.vgt_gs_onchip_cntl =
1725 S_028A44_ES_VERTS_PER_SUBGRP(shader->ngg.hw_max_esverts) |
1726 S_028A44_GS_PRIMS_PER_SUBGRP(shader->ngg.max_gsprims) |
1727 S_028A44_GS_INST_PRIMS_IN_SUBGRP(shader->ngg.max_gsprims * gs_num_invocations);
1728
1729 /* On gfx10, the GE only checks against the maximum number of ES verts after
1730 * allocating a full GS primitive. So we need to ensure that whenever
1731 * this check passes, there is enough space for a full primitive without
1732 * vertex reuse. VERT_GRP_SIZE=256 doesn't need this. We should always get 256
1733 * if we have enough LDS.
1734 *
1735 * Tessellation is unaffected because it always sets GE_CNTL.VERT_GRP_SIZE = 0.
1736 */
1737 if ((sscreen->info.gfx_level == GFX10) &&
1738 (es_stage == MESA_SHADER_VERTEX || gs_stage == MESA_SHADER_VERTEX) && /* = no tess */
1739 shader->ngg.hw_max_esverts != 256 &&
1740 shader->ngg.hw_max_esverts > 5) {
1741 /* This could be based on the input primitive type. 5 is the worst case
1742 * for primitive types with adjacency.
1743 */
1744 shader->ge_cntl &= C_03096C_VERT_GRP_SIZE;
1745 shader->ge_cntl |= S_03096C_VERT_GRP_SIZE(shader->ngg.hw_max_esverts - 5);
1746 }
1747 }
1748
1749 if (window_space) {
1750 shader->ngg.pa_cl_vte_cntl = S_028818_VTX_XY_FMT(1) | S_028818_VTX_Z_FMT(1);
1751 } else {
1752 shader->ngg.pa_cl_vte_cntl = S_028818_VTX_W0_FMT(1) |
1753 S_028818_VPORT_X_SCALE_ENA(1) | S_028818_VPORT_X_OFFSET_ENA(1) |
1754 S_028818_VPORT_Y_SCALE_ENA(1) | S_028818_VPORT_Y_OFFSET_ENA(1) |
1755 S_028818_VPORT_Z_SCALE_ENA(1) | S_028818_VPORT_Z_OFFSET_ENA(1);
1756 }
1757
1758 if (sscreen->info.gfx_level >= GFX12) {
1759 shader->ngg.vgt_shader_stages_en =
1760 S_028A98_GS_EN(gs_stage == MESA_SHADER_GEOMETRY) |
1761 S_028A98_PRIMGEN_PASSTHRU_NO_MSG(gfx10_is_ngg_passthrough(shader)) |
1762 S_028A98_GS_W32_EN(shader->wave_size == 32) |
1763 S_028A98_NGG_WAVE_ID_EN(si_shader_uses_streamout(shader));
1764 } else {
1765 shader->ngg.vgt_shader_stages_en =
1766 S_028B54_ES_EN(es_stage == MESA_SHADER_TESS_EVAL ?
1767 V_028B54_ES_STAGE_DS : V_028B54_ES_STAGE_REAL) |
1768 S_028B54_GS_EN(gs_stage == MESA_SHADER_GEOMETRY) |
1769 S_028B54_PRIMGEN_EN(1) |
1770 S_028B54_PRIMGEN_PASSTHRU_EN(gfx10_is_ngg_passthrough(shader)) |
1771 S_028B54_PRIMGEN_PASSTHRU_NO_MSG(gfx10_is_ngg_passthrough(shader) &&
1772 sscreen->info.family >= CHIP_NAVI23) |
1773 S_028B54_NGG_WAVE_ID_EN(si_shader_uses_streamout(shader)) |
1774 S_028B54_GS_W32_EN(shader->wave_size == 32) |
1775 S_028B54_MAX_PRIMGRP_IN_WAVE(2);
1776 }
1777
1778 ac_pm4_finalize(&pm4->base);
1779 }
1780
si_emit_shader_vs(struct si_context * sctx,unsigned index)1781 static void si_emit_shader_vs(struct si_context *sctx, unsigned index)
1782 {
1783 struct si_shader *shader = sctx->queued.named.vs;
1784
1785 radeon_begin(&sctx->gfx_cs);
1786 radeon_opt_set_context_reg(R_028A40_VGT_GS_MODE, SI_TRACKED_VGT_GS_MODE,
1787 shader->vs.vgt_gs_mode);
1788 radeon_opt_set_context_reg(R_028A84_VGT_PRIMITIVEID_EN, SI_TRACKED_VGT_PRIMITIVEID_EN,
1789 shader->vs.vgt_primitiveid_en);
1790
1791 if (sctx->gfx_level <= GFX8) {
1792 radeon_opt_set_context_reg(R_028AB4_VGT_REUSE_OFF, SI_TRACKED_VGT_REUSE_OFF,
1793 shader->vs.vgt_reuse_off);
1794 }
1795
1796 radeon_opt_set_context_reg(R_0286C4_SPI_VS_OUT_CONFIG, SI_TRACKED_SPI_VS_OUT_CONFIG,
1797 shader->vs.spi_vs_out_config);
1798
1799 radeon_opt_set_context_reg(R_02870C_SPI_SHADER_POS_FORMAT,
1800 SI_TRACKED_SPI_SHADER_POS_FORMAT,
1801 shader->vs.spi_shader_pos_format);
1802
1803 radeon_opt_set_context_reg(R_028818_PA_CL_VTE_CNTL, SI_TRACKED_PA_CL_VTE_CNTL,
1804 shader->vs.pa_cl_vte_cntl);
1805
1806 if (shader->selector->stage == MESA_SHADER_TESS_EVAL)
1807 radeon_opt_set_context_reg(R_028B6C_VGT_TF_PARAM, SI_TRACKED_VGT_TF_PARAM,
1808 shader->vgt_tf_param);
1809
1810 if (shader->vgt_vertex_reuse_block_cntl)
1811 radeon_opt_set_context_reg(R_028C58_VGT_VERTEX_REUSE_BLOCK_CNTL,
1812 SI_TRACKED_VGT_VERTEX_REUSE_BLOCK_CNTL,
1813 shader->vgt_vertex_reuse_block_cntl);
1814
1815 /* Required programming for tessellation. (legacy pipeline only) */
1816 if (sctx->gfx_level >= GFX10 && shader->selector->stage == MESA_SHADER_TESS_EVAL) {
1817 radeon_opt_set_context_reg(R_028A44_VGT_GS_ONCHIP_CNTL,
1818 SI_TRACKED_VGT_GS_ONCHIP_CNTL,
1819 S_028A44_ES_VERTS_PER_SUBGRP(250) |
1820 S_028A44_GS_PRIMS_PER_SUBGRP(126) |
1821 S_028A44_GS_INST_PRIMS_IN_SUBGRP(126));
1822 }
1823
1824 radeon_end_update_context_roll();
1825
1826 /* GE_PC_ALLOC is not a context register, so it doesn't cause a context roll. */
1827 if (sctx->gfx_level >= GFX10) {
1828 radeon_begin_again(&sctx->gfx_cs);
1829 radeon_opt_set_uconfig_reg(R_030980_GE_PC_ALLOC, SI_TRACKED_GE_PC_ALLOC,
1830 shader->vs.ge_pc_alloc);
1831 radeon_end();
1832 }
1833 }
1834
1835 /**
1836 * Compute the state for \p shader, which will run as a vertex shader on the
1837 * hardware.
1838 *
1839 * If \p gs is non-NULL, it points to the geometry shader for which this shader
1840 * is the copy shader.
1841 */
si_shader_vs(struct si_screen * sscreen,struct si_shader * shader,struct si_shader_selector * gs)1842 static void si_shader_vs(struct si_screen *sscreen, struct si_shader *shader,
1843 struct si_shader_selector *gs)
1844 {
1845 const struct si_shader_info *info = &shader->selector->info;
1846 struct si_pm4_state *pm4;
1847 unsigned num_user_sgprs, vgpr_comp_cnt;
1848 uint64_t va;
1849 unsigned nparams, oc_lds_en;
1850 bool window_space = shader->selector->stage == MESA_SHADER_VERTEX ?
1851 info->base.vs.window_space_position : 0;
1852 bool enable_prim_id = shader->key.ge.mono.u.vs_export_prim_id || info->uses_primid;
1853
1854 assert(sscreen->info.gfx_level < GFX11);
1855
1856 pm4 = si_get_shader_pm4_state(shader, si_emit_shader_vs);
1857 if (!pm4)
1858 return;
1859
1860 /* We always write VGT_GS_MODE in the VS state, because every switch
1861 * between different shader pipelines involving a different GS or no
1862 * GS at all involves a switch of the VS (different GS use different
1863 * copy shaders). On the other hand, when the API switches from a GS to
1864 * no GS and then back to the same GS used originally, the GS state is
1865 * not sent again.
1866 */
1867 if (!gs) {
1868 unsigned mode = V_028A40_GS_OFF;
1869
1870 /* PrimID needs GS scenario A. */
1871 if (enable_prim_id)
1872 mode = V_028A40_GS_SCENARIO_A;
1873
1874 shader->vs.vgt_gs_mode = S_028A40_MODE(mode);
1875 shader->vs.vgt_primitiveid_en = enable_prim_id;
1876 } else {
1877 shader->vs.vgt_gs_mode =
1878 ac_vgt_gs_mode(gs->info.base.gs.vertices_out, sscreen->info.gfx_level);
1879 shader->vs.vgt_primitiveid_en = 0;
1880 }
1881
1882 if (sscreen->info.gfx_level <= GFX8) {
1883 /* Reuse needs to be set off if we write oViewport. */
1884 shader->vs.vgt_reuse_off = S_028AB4_REUSE_OFF(info->writes_viewport_index);
1885 }
1886
1887 va = shader->bo->gpu_address;
1888
1889 if (gs) {
1890 vgpr_comp_cnt = 0; /* only VertexID is needed for GS-COPY. */
1891 num_user_sgprs = SI_GSCOPY_NUM_USER_SGPR;
1892 } else if (shader->selector->stage == MESA_SHADER_VERTEX) {
1893 vgpr_comp_cnt = si_get_vs_vgpr_comp_cnt(sscreen, shader, enable_prim_id);
1894
1895 if (info->base.vs.blit_sgprs_amd) {
1896 num_user_sgprs = SI_SGPR_VS_BLIT_DATA + info->base.vs.blit_sgprs_amd;
1897 } else {
1898 num_user_sgprs = si_get_num_vs_user_sgprs(shader, SI_VS_NUM_USER_SGPR);
1899 }
1900 } else if (shader->selector->stage == MESA_SHADER_TESS_EVAL) {
1901 vgpr_comp_cnt = enable_prim_id ? 3 : 2;
1902 num_user_sgprs = SI_TES_NUM_USER_SGPR;
1903 } else
1904 unreachable("invalid shader selector type");
1905
1906 /* VS is required to export at least one param. */
1907 nparams = MAX2(shader->info.nr_param_exports, 1);
1908 shader->vs.spi_vs_out_config = S_0286C4_VS_EXPORT_COUNT(nparams - 1);
1909
1910 if (sscreen->info.gfx_level >= GFX10) {
1911 shader->vs.spi_vs_out_config |=
1912 S_0286C4_NO_PC_EXPORT(shader->info.nr_param_exports == 0);
1913 }
1914
1915 shader->vs.spi_shader_pos_format =
1916 S_02870C_POS0_EXPORT_FORMAT(V_02870C_SPI_SHADER_4COMP) |
1917 S_02870C_POS1_EXPORT_FORMAT(shader->info.nr_pos_exports > 1 ? V_02870C_SPI_SHADER_4COMP
1918 : V_02870C_SPI_SHADER_NONE) |
1919 S_02870C_POS2_EXPORT_FORMAT(shader->info.nr_pos_exports > 2 ? V_02870C_SPI_SHADER_4COMP
1920 : V_02870C_SPI_SHADER_NONE) |
1921 S_02870C_POS3_EXPORT_FORMAT(shader->info.nr_pos_exports > 3 ? V_02870C_SPI_SHADER_4COMP
1922 : V_02870C_SPI_SHADER_NONE);
1923 unsigned late_alloc_wave64, cu_mask;
1924 ac_compute_late_alloc(&sscreen->info, false, false,
1925 shader->config.scratch_bytes_per_wave > 0,
1926 &late_alloc_wave64, &cu_mask);
1927
1928 shader->vs.ge_pc_alloc = S_030980_OVERSUB_EN(late_alloc_wave64 > 0) |
1929 S_030980_NUM_PC_LINES(sscreen->info.pc_lines / 4 - 1);
1930 shader->pa_cl_vs_out_cntl = si_get_vs_out_cntl(shader->selector, shader, false);
1931
1932 oc_lds_en = shader->selector->stage == MESA_SHADER_TESS_EVAL ? 1 : 0;
1933
1934 if (sscreen->info.gfx_level >= GFX7) {
1935 ac_pm4_set_reg_idx3(&pm4->base, R_00B118_SPI_SHADER_PGM_RSRC3_VS,
1936 ac_apply_cu_en(S_00B118_CU_EN(cu_mask) |
1937 S_00B118_WAVE_LIMIT(0x3F),
1938 C_00B118_CU_EN, 0, &sscreen->info));
1939 ac_pm4_set_reg(&pm4->base, R_00B11C_SPI_SHADER_LATE_ALLOC_VS, S_00B11C_LIMIT(late_alloc_wave64));
1940 }
1941
1942 ac_pm4_set_reg(&pm4->base, R_00B120_SPI_SHADER_PGM_LO_VS, va >> 8);
1943 ac_pm4_set_reg(&pm4->base, R_00B124_SPI_SHADER_PGM_HI_VS,
1944 S_00B124_MEM_BASE(sscreen->info.address32_hi >> 8));
1945
1946 uint32_t rsrc1 =
1947 S_00B128_VGPRS(si_shader_encode_vgprs(shader)) |
1948 S_00B128_SGPRS(si_shader_encode_sgprs(shader)) |
1949 S_00B128_VGPR_COMP_CNT(vgpr_comp_cnt) |
1950 S_00B128_DX10_CLAMP(1) |
1951 S_00B128_MEM_ORDERED(si_shader_mem_ordered(shader)) |
1952 S_00B128_FLOAT_MODE(shader->config.float_mode);
1953 uint32_t rsrc2 = S_00B12C_USER_SGPR(num_user_sgprs) | S_00B12C_OC_LDS_EN(oc_lds_en) |
1954 S_00B12C_SCRATCH_EN(shader->config.scratch_bytes_per_wave > 0);
1955
1956 if (sscreen->info.gfx_level >= GFX10)
1957 rsrc2 |= S_00B12C_USER_SGPR_MSB_GFX10(num_user_sgprs >> 5) |
1958 S_00B12C_SHARED_VGPR_CNT(shader->config.num_shared_vgprs / 8);
1959 else if (sscreen->info.gfx_level == GFX9)
1960 rsrc2 |= S_00B12C_USER_SGPR_MSB_GFX9(num_user_sgprs >> 5);
1961
1962 if (si_shader_uses_streamout(shader)) {
1963 rsrc2 |= S_00B12C_SO_BASE0_EN(!!shader->selector->info.base.xfb_stride[0]) |
1964 S_00B12C_SO_BASE1_EN(!!shader->selector->info.base.xfb_stride[1]) |
1965 S_00B12C_SO_BASE2_EN(!!shader->selector->info.base.xfb_stride[2]) |
1966 S_00B12C_SO_BASE3_EN(!!shader->selector->info.base.xfb_stride[3]) |
1967 S_00B12C_SO_EN(1);
1968 }
1969
1970 ac_pm4_set_reg(&pm4->base, R_00B128_SPI_SHADER_PGM_RSRC1_VS, rsrc1);
1971 ac_pm4_set_reg(&pm4->base, R_00B12C_SPI_SHADER_PGM_RSRC2_VS, rsrc2);
1972
1973 if (window_space)
1974 shader->vs.pa_cl_vte_cntl = S_028818_VTX_XY_FMT(1) | S_028818_VTX_Z_FMT(1);
1975 else
1976 shader->vs.pa_cl_vte_cntl =
1977 S_028818_VTX_W0_FMT(1) | S_028818_VPORT_X_SCALE_ENA(1) | S_028818_VPORT_X_OFFSET_ENA(1) |
1978 S_028818_VPORT_Y_SCALE_ENA(1) | S_028818_VPORT_Y_OFFSET_ENA(1) |
1979 S_028818_VPORT_Z_SCALE_ENA(1) | S_028818_VPORT_Z_OFFSET_ENA(1);
1980
1981 if (shader->selector->stage == MESA_SHADER_TESS_EVAL)
1982 si_set_tesseval_regs(sscreen, shader->selector, shader);
1983
1984 polaris_set_vgt_vertex_reuse(sscreen, shader->selector, shader);
1985 ac_pm4_finalize(&pm4->base);
1986 }
1987
si_get_spi_shader_col_format(struct si_shader * shader)1988 static unsigned si_get_spi_shader_col_format(struct si_shader *shader)
1989 {
1990 unsigned spi_shader_col_format = shader->key.ps.part.epilog.spi_shader_col_format;
1991 unsigned value = 0, num_mrts = 0;
1992 unsigned i, num_targets = (util_last_bit(spi_shader_col_format) + 3) / 4;
1993
1994 /* Remove holes in spi_shader_col_format. */
1995 for (i = 0; i < num_targets; i++) {
1996 unsigned spi_format = (spi_shader_col_format >> (i * 4)) & 0xf;
1997
1998 if (spi_format) {
1999 value |= spi_format << (num_mrts * 4);
2000 num_mrts++;
2001 }
2002 }
2003
2004 return value;
2005 }
2006
gfx6_emit_shader_ps(struct si_context * sctx,unsigned index)2007 static void gfx6_emit_shader_ps(struct si_context *sctx, unsigned index)
2008 {
2009 struct si_shader *shader = sctx->queued.named.ps;
2010
2011 radeon_begin(&sctx->gfx_cs);
2012 radeon_opt_set_context_reg2(R_0286CC_SPI_PS_INPUT_ENA, SI_TRACKED_SPI_PS_INPUT_ENA,
2013 shader->ps.spi_ps_input_ena,
2014 shader->ps.spi_ps_input_addr);
2015 radeon_opt_set_context_reg(R_0286D8_SPI_PS_IN_CONTROL, SI_TRACKED_SPI_PS_IN_CONTROL,
2016 shader->ps.spi_ps_in_control);
2017 radeon_opt_set_context_reg2(R_028710_SPI_SHADER_Z_FORMAT, SI_TRACKED_SPI_SHADER_Z_FORMAT,
2018 shader->ps.spi_shader_z_format,
2019 shader->ps.spi_shader_col_format);
2020 radeon_opt_set_context_reg(R_02823C_CB_SHADER_MASK, SI_TRACKED_CB_SHADER_MASK,
2021 shader->ps.cb_shader_mask);
2022 radeon_end_update_context_roll();
2023 }
2024
gfx11_dgpu_emit_shader_ps(struct si_context * sctx,unsigned index)2025 static void gfx11_dgpu_emit_shader_ps(struct si_context *sctx, unsigned index)
2026 {
2027 struct si_shader *shader = sctx->queued.named.ps;
2028
2029 radeon_begin(&sctx->gfx_cs);
2030 gfx11_begin_packed_context_regs();
2031 gfx11_opt_set_context_reg(R_0286CC_SPI_PS_INPUT_ENA, SI_TRACKED_SPI_PS_INPUT_ENA,
2032 shader->ps.spi_ps_input_ena);
2033 gfx11_opt_set_context_reg(R_0286D0_SPI_PS_INPUT_ADDR, SI_TRACKED_SPI_PS_INPUT_ADDR,
2034 shader->ps.spi_ps_input_addr);
2035 gfx11_opt_set_context_reg(R_0286D8_SPI_PS_IN_CONTROL, SI_TRACKED_SPI_PS_IN_CONTROL,
2036 shader->ps.spi_ps_in_control);
2037 gfx11_opt_set_context_reg(R_028710_SPI_SHADER_Z_FORMAT, SI_TRACKED_SPI_SHADER_Z_FORMAT,
2038 shader->ps.spi_shader_z_format);
2039 gfx11_opt_set_context_reg(R_028714_SPI_SHADER_COL_FORMAT, SI_TRACKED_SPI_SHADER_COL_FORMAT,
2040 shader->ps.spi_shader_col_format);
2041 gfx11_opt_set_context_reg(R_02823C_CB_SHADER_MASK, SI_TRACKED_CB_SHADER_MASK,
2042 shader->ps.cb_shader_mask);
2043 gfx11_end_packed_context_regs();
2044 radeon_end(); /* don't track context rolls on GFX11 */
2045 }
2046
gfx12_emit_shader_ps(struct si_context * sctx,unsigned index)2047 static void gfx12_emit_shader_ps(struct si_context *sctx, unsigned index)
2048 {
2049 struct si_shader *shader = sctx->queued.named.ps;
2050
2051 radeon_begin(&sctx->gfx_cs);
2052 gfx12_begin_context_regs();
2053 gfx12_opt_set_context_reg(R_028640_SPI_PS_IN_CONTROL, SI_TRACKED_SPI_PS_IN_CONTROL,
2054 shader->ps.spi_ps_in_control);
2055 gfx12_opt_set_context_reg(R_028650_SPI_SHADER_Z_FORMAT, SI_TRACKED_SPI_SHADER_Z_FORMAT,
2056 shader->ps.spi_shader_z_format);
2057 gfx12_opt_set_context_reg(R_028654_SPI_SHADER_COL_FORMAT, SI_TRACKED_SPI_SHADER_COL_FORMAT,
2058 shader->ps.spi_shader_col_format);
2059 gfx12_opt_set_context_reg(R_02865C_SPI_PS_INPUT_ENA, SI_TRACKED_SPI_PS_INPUT_ENA,
2060 shader->ps.spi_ps_input_ena);
2061 gfx12_opt_set_context_reg(R_028660_SPI_PS_INPUT_ADDR, SI_TRACKED_SPI_PS_INPUT_ADDR,
2062 shader->ps.spi_ps_input_addr);
2063 gfx12_opt_set_context_reg(R_028854_CB_SHADER_MASK, SI_TRACKED_CB_SHADER_MASK,
2064 shader->ps.cb_shader_mask);
2065 gfx12_opt_set_context_reg(R_028BBC_PA_SC_HISZ_CONTROL, SI_TRACKED_PA_SC_HISZ_CONTROL,
2066 shader->ps.pa_sc_hisz_control);
2067 gfx12_end_context_regs();
2068 radeon_end(); /* don't track context rolls on GFX12 */
2069 }
2070
si_shader_ps(struct si_screen * sscreen,struct si_shader * shader)2071 static void si_shader_ps(struct si_screen *sscreen, struct si_shader *shader)
2072 {
2073 struct si_shader_info *info = &shader->selector->info;
2074 const unsigned input_ena = shader->config.spi_ps_input_ena;
2075 /* At least one of these is required to be set. */
2076 ASSERTED unsigned num_required_vgpr_inputs =
2077 G_0286CC_PERSP_SAMPLE_ENA(input_ena) + G_0286CC_PERSP_CENTER_ENA(input_ena) +
2078 G_0286CC_PERSP_CENTROID_ENA(input_ena) + G_0286CC_PERSP_PULL_MODEL_ENA(input_ena) +
2079 G_0286CC_LINEAR_SAMPLE_ENA(input_ena) + G_0286CC_LINEAR_CENTER_ENA(input_ena) +
2080 G_0286CC_LINEAR_CENTROID_ENA(input_ena) + G_0286CC_LINE_STIPPLE_TEX_ENA(input_ena);
2081
2082 /* we need to enable at least one of them, otherwise we hang the GPU */
2083 assert(num_required_vgpr_inputs > 0);
2084 /* POS_W_FLOAT_ENA requires one of the perspective weights. */
2085 assert(!G_0286CC_POS_W_FLOAT_ENA(input_ena) || G_0286CC_PERSP_SAMPLE_ENA(input_ena) ||
2086 G_0286CC_PERSP_CENTER_ENA(input_ena) || G_0286CC_PERSP_CENTROID_ENA(input_ena) ||
2087 G_0286CC_PERSP_PULL_MODEL_ENA(input_ena));
2088
2089 /* Validate interpolation optimization flags (read as implications). */
2090 assert(!shader->key.ps.part.prolog.bc_optimize_for_persp ||
2091 (G_0286CC_PERSP_CENTER_ENA(input_ena) && G_0286CC_PERSP_CENTROID_ENA(input_ena)));
2092 assert(!shader->key.ps.part.prolog.bc_optimize_for_linear ||
2093 (G_0286CC_LINEAR_CENTER_ENA(input_ena) && G_0286CC_LINEAR_CENTROID_ENA(input_ena)));
2094 assert(!shader->key.ps.part.prolog.force_persp_center_interp || num_required_vgpr_inputs == 1 ||
2095 (!G_0286CC_PERSP_SAMPLE_ENA(input_ena) && !G_0286CC_PERSP_CENTROID_ENA(input_ena)));
2096 assert(!shader->key.ps.part.prolog.force_linear_center_interp || num_required_vgpr_inputs == 1 ||
2097 (!G_0286CC_LINEAR_SAMPLE_ENA(input_ena) && !G_0286CC_LINEAR_CENTROID_ENA(input_ena)));
2098 assert(!shader->key.ps.part.prolog.force_persp_sample_interp || num_required_vgpr_inputs == 1 ||
2099 (!G_0286CC_PERSP_CENTER_ENA(input_ena) && !G_0286CC_PERSP_CENTROID_ENA(input_ena)));
2100 assert(!shader->key.ps.part.prolog.force_linear_sample_interp || num_required_vgpr_inputs == 1 ||
2101 (!G_0286CC_LINEAR_CENTER_ENA(input_ena) && !G_0286CC_LINEAR_CENTROID_ENA(input_ena)));
2102
2103 /* color_two_side always enables FRONT_FACE. Since st/mesa disables two-side colors if the back
2104 * face is culled, the only case when both color_two_side and force_front_face_input can be set
2105 * is when the front face is culled (which means force_front_face_input == -1).
2106 */
2107 assert(!shader->key.ps.opt.force_front_face_input || !G_0286CC_FRONT_FACE_ENA(input_ena) ||
2108 (shader->key.ps.part.prolog.color_two_side &&
2109 shader->key.ps.opt.force_front_face_input == -1));
2110
2111 /* Validate cases when the optimizations are off (read as implications). */
2112 assert(shader->key.ps.part.prolog.bc_optimize_for_persp ||
2113 !G_0286CC_PERSP_CENTER_ENA(input_ena) || !G_0286CC_PERSP_CENTROID_ENA(input_ena));
2114 assert(shader->key.ps.part.prolog.bc_optimize_for_linear ||
2115 !G_0286CC_LINEAR_CENTER_ENA(input_ena) || !G_0286CC_LINEAR_CENTROID_ENA(input_ena));
2116
2117 /* DB_SHADER_CONTROL */
2118 shader->ps.db_shader_control =
2119 S_02880C_Z_EXPORT_ENABLE(shader->ps.writes_z) |
2120 S_02880C_STENCIL_TEST_VAL_EXPORT_ENABLE(shader->ps.writes_stencil) |
2121 S_02880C_MASK_EXPORT_ENABLE(shader->ps.writes_samplemask) |
2122 S_02880C_COVERAGE_TO_MASK_ENABLE(sscreen->info.gfx_level <= GFX10_3 &&
2123 shader->key.ps.part.epilog.alpha_to_coverage_via_mrtz) |
2124 S_02880C_KILL_ENABLE(si_shader_uses_discard(shader));
2125
2126 if (sscreen->info.gfx_level >= GFX12)
2127 shader->ps.pa_sc_hisz_control = S_028BBC_ROUND(2); /* required minimum value */
2128
2129 switch (info->base.fs.depth_layout) {
2130 case FRAG_DEPTH_LAYOUT_GREATER:
2131 shader->ps.db_shader_control |= S_02880C_CONSERVATIVE_Z_EXPORT(V_02880C_EXPORT_GREATER_THAN_Z);
2132 if (sscreen->info.gfx_level >= GFX12)
2133 shader->ps.pa_sc_hisz_control |= S_028BBC_CONSERVATIVE_Z_EXPORT(V_028BBC_EXPORT_GREATER_THAN_Z);
2134 break;
2135 case FRAG_DEPTH_LAYOUT_LESS:
2136 shader->ps.db_shader_control |= S_02880C_CONSERVATIVE_Z_EXPORT(V_02880C_EXPORT_LESS_THAN_Z);
2137 if (sscreen->info.gfx_level >= GFX12)
2138 shader->ps.pa_sc_hisz_control |= S_028BBC_CONSERVATIVE_Z_EXPORT(V_028BBC_EXPORT_LESS_THAN_Z);
2139 break;
2140 default:;
2141 }
2142
2143 /* Z_ORDER, EXEC_ON_HIER_FAIL and EXEC_ON_NOOP should be set as following:
2144 *
2145 * | early Z/S | writes_mem | allow_ReZ? | Z_ORDER | EXEC_ON_HIER_FAIL | EXEC_ON_NOOP
2146 * --|-----------|------------|------------|--------------------|-------------------|-------------
2147 * 1a| false | false | true | EarlyZ_Then_ReZ | 0 | 0
2148 * 1b| false | false | false | EarlyZ_Then_LateZ | 0 | 0
2149 * 2 | false | true | n/a | LateZ | 1 | 0
2150 * 3 | true | false | n/a | EarlyZ_Then_LateZ | 0 | 0
2151 * 4 | true | true | n/a | EarlyZ_Then_LateZ | 0 | 1
2152 *
2153 * In cases 3 and 4, HW will force Z_ORDER to EarlyZ regardless of what's set in the register.
2154 * In case 2, NOOP_CULL is a don't care field. In case 2, 3 and 4, ReZ doesn't make sense.
2155 *
2156 * Don't use ReZ without profiling !!!
2157 *
2158 * ReZ decreases performance by 15% in DiRT: Showdown on Ultra settings, which has pretty complex
2159 * shaders.
2160 */
2161 if (info->base.fs.early_fragment_tests) {
2162 /* Cases 3, 4. */
2163 shader->ps.db_shader_control |= S_02880C_DEPTH_BEFORE_SHADER(1) |
2164 S_02880C_Z_ORDER(V_02880C_EARLY_Z_THEN_LATE_Z) |
2165 S_02880C_EXEC_ON_NOOP(info->base.writes_memory);
2166 } else if (info->base.writes_memory) {
2167 /* Case 2. */
2168 shader->ps.db_shader_control |= S_02880C_Z_ORDER(V_02880C_LATE_Z) |
2169 S_02880C_EXEC_ON_HIER_FAIL(1);
2170 } else {
2171 /* Case 1. */
2172 shader->ps.db_shader_control |= S_02880C_Z_ORDER(V_02880C_EARLY_Z_THEN_LATE_Z);
2173 }
2174
2175 if (info->base.fs.post_depth_coverage)
2176 shader->ps.db_shader_control |= S_02880C_PRE_SHADER_DEPTH_COVERAGE_ENABLE(1);
2177
2178 /* Bug workaround for smoothing (overrasterization) on GFX6. */
2179 if (sscreen->info.gfx_level == GFX6 && shader->key.ps.mono.poly_line_smoothing) {
2180 shader->ps.db_shader_control &= C_02880C_Z_ORDER;
2181 shader->ps.db_shader_control |= S_02880C_Z_ORDER(V_02880C_LATE_Z);
2182 }
2183
2184 if (sscreen->info.has_rbplus && !sscreen->info.rbplus_allowed)
2185 shader->ps.db_shader_control |= S_02880C_DUAL_QUAD_DISABLE(1);
2186
2187 shader->ps.spi_shader_col_format = si_get_spi_shader_col_format(shader);
2188 shader->ps.cb_shader_mask = ac_get_cb_shader_mask(shader->key.ps.part.epilog.spi_shader_col_format);
2189 shader->ps.spi_ps_input_ena = shader->config.spi_ps_input_ena;
2190 shader->ps.spi_ps_input_addr = shader->config.spi_ps_input_addr;
2191 shader->ps.num_interp = si_get_ps_num_interp(shader);
2192 shader->ps.spi_shader_z_format =
2193 ac_get_spi_shader_z_format(shader->ps.writes_z, shader->ps.writes_stencil,
2194 shader->ps.writes_samplemask,
2195 shader->key.ps.part.epilog.alpha_to_coverage_via_mrtz);
2196
2197 /* Ensure that some export memory is always allocated, for two reasons:
2198 *
2199 * 1) Correctness: The hardware ignores the EXEC mask if no export
2200 * memory is allocated, so KILL and alpha test do not work correctly
2201 * without this.
2202 * 2) Performance: Every shader needs at least a NULL export, even when
2203 * it writes no color/depth output. The NULL export instruction
2204 * stalls without this setting.
2205 *
2206 * Don't add this to CB_SHADER_MASK.
2207 *
2208 * GFX10 supports pixel shaders without exports by setting both
2209 * the color and Z formats to SPI_SHADER_ZERO. The hw will skip export
2210 * instructions if any are present.
2211 *
2212 * RB+ depth-only rendering requires SPI_SHADER_32_R.
2213 */
2214 bool has_mrtz = shader->ps.spi_shader_z_format != V_028710_SPI_SHADER_ZERO;
2215
2216 if (!shader->ps.spi_shader_col_format) {
2217 if (shader->key.ps.part.epilog.rbplus_depth_only_opt) {
2218 shader->ps.spi_shader_col_format = V_028714_SPI_SHADER_32_R;
2219 } else if (!has_mrtz) {
2220 if (sscreen->info.gfx_level >= GFX10) {
2221 if (G_02880C_KILL_ENABLE(shader->ps.db_shader_control))
2222 shader->ps.spi_shader_col_format = V_028714_SPI_SHADER_32_R;
2223 } else {
2224 shader->ps.spi_shader_col_format = V_028714_SPI_SHADER_32_R;
2225 }
2226 }
2227 }
2228
2229 if (sscreen->info.gfx_level >= GFX12) {
2230 shader->ps.spi_ps_in_control = S_028640_PARAM_GEN(shader->key.ps.mono.point_smoothing) |
2231 S_028640_PS_W32_EN(shader->wave_size == 32);
2232 shader->ps.spi_gs_out_config_ps = S_00B0C4_NUM_INTERP(shader->ps.num_interp);
2233 } else {
2234 /* Enable PARAM_GEN for point smoothing.
2235 * Gfx11 workaround when there are no PS inputs but LDS is used.
2236 */
2237 bool param_gen = shader->key.ps.mono.point_smoothing ||
2238 (sscreen->info.gfx_level == GFX11 && !shader->ps.num_interp &&
2239 shader->config.lds_size);
2240
2241 shader->ps.spi_ps_in_control = S_0286D8_NUM_INTERP(shader->ps.num_interp) |
2242 S_0286D8_PARAM_GEN(param_gen) |
2243 S_0286D8_PS_W32_EN(shader->wave_size == 32);
2244 }
2245
2246 struct si_pm4_state *pm4 = si_get_shader_pm4_state(shader, NULL);
2247 if (!pm4)
2248 return;
2249
2250 if (sscreen->info.gfx_level >= GFX12)
2251 pm4->atom.emit = gfx12_emit_shader_ps;
2252 else if (sscreen->info.has_set_context_pairs_packed)
2253 pm4->atom.emit = gfx11_dgpu_emit_shader_ps;
2254 else
2255 pm4->atom.emit = gfx6_emit_shader_ps;
2256
2257 /* If multiple state sets are allowed to be in a bin, break the batch on a new PS. */
2258 if (sscreen->dpbb_allowed &&
2259 (sscreen->pbb_context_states_per_bin > 1 ||
2260 sscreen->pbb_persistent_states_per_bin > 1)) {
2261 ac_pm4_cmd_add(&pm4->base, PKT3(PKT3_EVENT_WRITE, 0, 0));
2262 ac_pm4_cmd_add(&pm4->base, EVENT_TYPE(V_028A90_BREAK_BATCH) | EVENT_INDEX(0));
2263 }
2264
2265 if (sscreen->info.gfx_level >= GFX12) {
2266 ac_pm4_set_reg(&pm4->base, R_00B01C_SPI_SHADER_PGM_RSRC4_PS,
2267 S_00B01C_WAVE_LIMIT_GFX12(0x3FF) |
2268 S_00B01C_LDS_GROUP_SIZE_GFX12(1) |
2269 S_00B01C_INST_PREF_SIZE(si_get_shader_prefetch_size(shader)));
2270 } else if (sscreen->info.gfx_level >= GFX11) {
2271 unsigned cu_mask_ps = ac_gfx103_get_cu_mask_ps(&sscreen->info);
2272
2273 ac_pm4_set_reg_idx3(&pm4->base, R_00B004_SPI_SHADER_PGM_RSRC4_PS,
2274 ac_apply_cu_en(S_00B004_CU_EN(cu_mask_ps >> 16) |
2275 S_00B004_INST_PREF_SIZE(si_get_shader_prefetch_size(shader)),
2276 C_00B004_CU_EN, 16, &sscreen->info));
2277 }
2278
2279 uint64_t va = shader->bo->gpu_address;
2280 ac_pm4_set_reg(&pm4->base, R_00B020_SPI_SHADER_PGM_LO_PS, va >> 8);
2281 ac_pm4_set_reg(&pm4->base, R_00B024_SPI_SHADER_PGM_HI_PS,
2282 S_00B024_MEM_BASE(sscreen->info.address32_hi >> 8));
2283
2284 ac_pm4_set_reg(&pm4->base, R_00B028_SPI_SHADER_PGM_RSRC1_PS,
2285 S_00B028_VGPRS(si_shader_encode_vgprs(shader)) |
2286 S_00B028_SGPRS(si_shader_encode_sgprs(shader)) |
2287 S_00B028_DX10_CLAMP(sscreen->info.gfx_level < GFX12) |
2288 S_00B028_MEM_ORDERED(si_shader_mem_ordered(shader)) |
2289 S_00B028_FLOAT_MODE(shader->config.float_mode));
2290 ac_pm4_set_reg(&pm4->base, R_00B02C_SPI_SHADER_PGM_RSRC2_PS,
2291 S_00B02C_EXTRA_LDS_SIZE(shader->config.lds_size) |
2292 S_00B02C_USER_SGPR(SI_PS_NUM_USER_SGPR) |
2293 S_00B02C_SCRATCH_EN(shader->config.scratch_bytes_per_wave > 0) |
2294 S_00B02C_SHARED_VGPR_CNT(shader->config.num_shared_vgprs / 8));
2295 ac_pm4_finalize(&pm4->base);
2296 }
2297
si_shader_init_pm4_state(struct si_screen * sscreen,struct si_shader * shader)2298 static void si_shader_init_pm4_state(struct si_screen *sscreen, struct si_shader *shader)
2299 {
2300 assert(shader->wave_size);
2301
2302 switch (shader->selector->stage) {
2303 case MESA_SHADER_VERTEX:
2304 if (shader->key.ge.as_ls)
2305 si_shader_ls(sscreen, shader);
2306 else if (shader->key.ge.as_es)
2307 si_shader_es(sscreen, shader);
2308 else if (shader->key.ge.as_ngg)
2309 gfx10_shader_ngg(sscreen, shader);
2310 else
2311 si_shader_vs(sscreen, shader, NULL);
2312 break;
2313 case MESA_SHADER_TESS_CTRL:
2314 si_shader_hs(sscreen, shader);
2315 break;
2316 case MESA_SHADER_TESS_EVAL:
2317 if (shader->key.ge.as_es)
2318 si_shader_es(sscreen, shader);
2319 else if (shader->key.ge.as_ngg)
2320 gfx10_shader_ngg(sscreen, shader);
2321 else
2322 si_shader_vs(sscreen, shader, NULL);
2323 break;
2324 case MESA_SHADER_GEOMETRY:
2325 if (shader->key.ge.as_ngg) {
2326 gfx10_shader_ngg(sscreen, shader);
2327 } else {
2328 /* VS must be initialized first because GS uses its fields. */
2329 si_shader_vs(sscreen, shader->gs_copy_shader, shader->selector);
2330 si_shader_gs(sscreen, shader);
2331 }
2332 break;
2333 case MESA_SHADER_FRAGMENT:
2334 si_shader_ps(sscreen, shader);
2335 break;
2336 default:
2337 assert(0);
2338 }
2339
2340 assert(!(sscreen->debug_flags & DBG(SQTT)) || shader->pm4.base.spi_shader_pgm_lo_reg != 0);
2341 }
2342
si_clear_vs_key_inputs(union si_shader_key * key)2343 static void si_clear_vs_key_inputs(union si_shader_key *key)
2344 {
2345 key->ge.mono.instance_divisor_is_one = 0;
2346 key->ge.mono.instance_divisor_is_fetched = 0;
2347 key->ge.mono.vs_fetch_opencode = 0;
2348 memset(key->ge.mono.vs_fix_fetch, 0, sizeof(key->ge.mono.vs_fix_fetch));
2349 }
2350
si_vs_key_update_inputs(struct si_context * sctx)2351 void si_vs_key_update_inputs(struct si_context *sctx)
2352 {
2353 struct si_shader_selector *vs = sctx->shader.vs.cso;
2354 struct si_vertex_elements *elts = sctx->vertex_elements;
2355 union si_shader_key *key = &sctx->shader.vs.key;
2356
2357 if (!vs)
2358 return;
2359
2360 if (vs->info.base.vs.blit_sgprs_amd) {
2361 si_clear_vs_key_inputs(key);
2362 key->ge.opt.prefer_mono = 0;
2363 sctx->uses_nontrivial_vs_inputs = false;
2364 return;
2365 }
2366
2367 bool uses_nontrivial_vs_inputs = false;
2368
2369 if (elts->instance_divisor_is_one || elts->instance_divisor_is_fetched)
2370 uses_nontrivial_vs_inputs = true;
2371
2372 key->ge.mono.instance_divisor_is_one = elts->instance_divisor_is_one;
2373 key->ge.mono.instance_divisor_is_fetched = elts->instance_divisor_is_fetched;
2374 key->ge.opt.prefer_mono = elts->instance_divisor_is_fetched;
2375
2376 unsigned count_mask = (1 << vs->info.num_inputs) - 1;
2377 unsigned fix = elts->fix_fetch_always & count_mask;
2378 unsigned opencode = elts->fix_fetch_opencode & count_mask;
2379
2380 if (sctx->vertex_buffer_unaligned & elts->vb_alignment_check_mask) {
2381 uint32_t mask = elts->fix_fetch_unaligned & count_mask;
2382 while (mask) {
2383 unsigned i = u_bit_scan(&mask);
2384 unsigned log_hw_load_size = 1 + ((elts->hw_load_is_dword >> i) & 1);
2385 unsigned vbidx = elts->vertex_buffer_index[i];
2386 const struct pipe_vertex_buffer *vb = &sctx->vertex_buffer[vbidx];
2387 unsigned align_mask = (1 << log_hw_load_size) - 1;
2388 if (vb->buffer_offset & align_mask) {
2389 fix |= 1 << i;
2390 opencode |= 1 << i;
2391 }
2392 }
2393 }
2394
2395 memset(key->ge.mono.vs_fix_fetch, 0, sizeof(key->ge.mono.vs_fix_fetch));
2396
2397 while (fix) {
2398 unsigned i = u_bit_scan(&fix);
2399 uint8_t fix_fetch = elts->fix_fetch[i];
2400
2401 key->ge.mono.vs_fix_fetch[i].bits = fix_fetch;
2402 if (fix_fetch)
2403 uses_nontrivial_vs_inputs = true;
2404 }
2405 key->ge.mono.vs_fetch_opencode = opencode;
2406 if (opencode)
2407 uses_nontrivial_vs_inputs = true;
2408
2409 sctx->uses_nontrivial_vs_inputs = uses_nontrivial_vs_inputs;
2410
2411 /* draw_vertex_state (display lists) requires that all VS input lowering is disabled
2412 * because its vertex elements never need any lowering.
2413 *
2414 * We just computed the key because we needed to set uses_nontrivial_vs_inputs, so that we know
2415 * whether the VS should be updated when we switch from draw_vertex_state to draw_vbo. Now
2416 * clear the VS input bits for draw_vertex_state. This should happen rarely because VS inputs
2417 * don't usually need any lowering.
2418 */
2419 if (uses_nontrivial_vs_inputs && sctx->force_trivial_vs_inputs)
2420 si_clear_vs_key_inputs(key);
2421 }
2422
si_get_vs_key_inputs(struct si_context * sctx,union si_shader_key * key)2423 static void si_get_vs_key_inputs(struct si_context *sctx, union si_shader_key *key)
2424 {
2425 key->ge.mono.instance_divisor_is_one = sctx->shader.vs.key.ge.mono.instance_divisor_is_one;
2426 key->ge.mono.instance_divisor_is_fetched = sctx->shader.vs.key.ge.mono.instance_divisor_is_fetched;
2427 key->ge.mono.vs_fetch_opencode = sctx->shader.vs.key.ge.mono.vs_fetch_opencode;
2428 memcpy(key->ge.mono.vs_fix_fetch, sctx->shader.vs.key.ge.mono.vs_fix_fetch,
2429 sizeof(key->ge.mono.vs_fix_fetch));
2430 }
2431
si_update_ps_inputs_read_or_disabled(struct si_context * sctx)2432 void si_update_ps_inputs_read_or_disabled(struct si_context *sctx)
2433 {
2434 struct si_shader_selector *ps = sctx->shader.ps.cso;
2435
2436 /* Find out if PS is disabled. */
2437 bool ps_disabled = true;
2438 if (ps) {
2439 bool ps_modifies_zs = ps->info.base.fs.uses_discard ||
2440 ps->info.writes_z ||
2441 ps->info.writes_stencil ||
2442 ps->info.writes_samplemask ||
2443 sctx->queued.named.blend->alpha_to_coverage ||
2444 sctx->queued.named.dsa->alpha_func != PIPE_FUNC_ALWAYS ||
2445 sctx->queued.named.rasterizer->poly_stipple_enable ||
2446 sctx->queued.named.rasterizer->point_smooth;
2447
2448 ps_disabled = sctx->queued.named.rasterizer->rasterizer_discard ||
2449 (!ps_modifies_zs && !ps->info.base.writes_memory &&
2450 !si_any_colorbuffer_written(sctx));
2451 }
2452
2453 uint64_t ps_inputs_read_or_disabled;
2454
2455 if (ps_disabled) {
2456 ps_inputs_read_or_disabled = 0;
2457 } else {
2458 uint64_t inputs_read = ps->info.inputs_read;
2459
2460 if (ps->info.colors_read && sctx->queued.named.rasterizer->two_side) {
2461 if (inputs_read & BITFIELD64_BIT(SI_UNIQUE_SLOT_COL0))
2462 inputs_read |= BITFIELD64_BIT(SI_UNIQUE_SLOT_BFC0);
2463
2464 if (inputs_read & BITFIELD64_BIT(SI_UNIQUE_SLOT_COL1))
2465 inputs_read |= BITFIELD64_BIT(SI_UNIQUE_SLOT_BFC1);
2466 }
2467
2468 ps_inputs_read_or_disabled = inputs_read;
2469 }
2470
2471 if (sctx->ps_inputs_read_or_disabled != ps_inputs_read_or_disabled) {
2472 sctx->ps_inputs_read_or_disabled = ps_inputs_read_or_disabled;
2473 sctx->do_update_shaders = true;
2474 }
2475 }
2476
si_vs_ps_key_update_rast_prim_smooth_stipple(struct si_context * sctx)2477 void si_vs_ps_key_update_rast_prim_smooth_stipple(struct si_context *sctx)
2478 {
2479 struct si_shader_ctx_state *hw_vs = si_get_vs(sctx);
2480 struct si_shader_selector *ps = sctx->shader.ps.cso;
2481
2482 if (!hw_vs->cso || !ps)
2483 return;
2484
2485 struct si_state_rasterizer *rs = sctx->queued.named.rasterizer;
2486 union si_shader_key *vs_key = &hw_vs->key; /* could also be TES or GS before PS */
2487 union si_shader_key *ps_key = &sctx->shader.ps.key;
2488
2489 bool old_kill_pointsize = vs_key->ge.opt.kill_pointsize;
2490 bool old_color_two_side = ps_key->ps.part.prolog.color_two_side;
2491 bool old_poly_stipple = ps_key->ps.part.prolog.poly_stipple;
2492 bool old_poly_line_smoothing = ps_key->ps.mono.poly_line_smoothing;
2493 bool old_point_smoothing = ps_key->ps.mono.point_smoothing;
2494 int old_force_front_face_input = ps_key->ps.opt.force_front_face_input;
2495
2496 if (sctx->current_rast_prim == MESA_PRIM_POINTS) {
2497 vs_key->ge.opt.kill_pointsize = 0;
2498 ps_key->ps.part.prolog.color_two_side = 0;
2499 ps_key->ps.part.prolog.poly_stipple = 0;
2500 ps_key->ps.mono.poly_line_smoothing = 0;
2501 ps_key->ps.mono.point_smoothing = rs->point_smooth;
2502 ps_key->ps.opt.force_front_face_input = ps->info.uses_frontface;
2503 } else if (util_prim_is_lines(sctx->current_rast_prim)) {
2504 vs_key->ge.opt.kill_pointsize = hw_vs->cso->info.writes_psize;
2505 ps_key->ps.part.prolog.color_two_side = 0;
2506 ps_key->ps.part.prolog.poly_stipple = 0;
2507 ps_key->ps.mono.poly_line_smoothing = rs->line_smooth && sctx->framebuffer.nr_samples <= 1;
2508 ps_key->ps.mono.point_smoothing = 0;
2509 ps_key->ps.opt.force_front_face_input = ps->info.uses_frontface;
2510 } else {
2511 /* Triangles. */
2512 vs_key->ge.opt.kill_pointsize = hw_vs->cso->info.writes_psize &&
2513 !rs->polygon_mode_is_points;
2514 ps_key->ps.part.prolog.color_two_side = rs->two_side && ps->info.colors_read;
2515 ps_key->ps.part.prolog.poly_stipple = rs->poly_stipple_enable;
2516 ps_key->ps.mono.poly_line_smoothing = rs->poly_smooth && sctx->framebuffer.nr_samples <= 1;
2517 ps_key->ps.mono.point_smoothing = 0;
2518 ps_key->ps.opt.force_front_face_input = ps->info.uses_frontface ? rs->force_front_face_input : 0;
2519 }
2520
2521 if (vs_key->ge.opt.kill_pointsize != old_kill_pointsize ||
2522 ps_key->ps.part.prolog.color_two_side != old_color_two_side ||
2523 ps_key->ps.part.prolog.poly_stipple != old_poly_stipple ||
2524 ps_key->ps.mono.poly_line_smoothing != old_poly_line_smoothing ||
2525 ps_key->ps.mono.point_smoothing != old_point_smoothing ||
2526 ps_key->ps.opt.force_front_face_input != old_force_front_face_input)
2527 sctx->do_update_shaders = true;
2528 }
2529
si_get_vs_key_outputs(struct si_context * sctx,struct si_shader_selector * vs,union si_shader_key * key)2530 static void si_get_vs_key_outputs(struct si_context *sctx, struct si_shader_selector *vs,
2531 union si_shader_key *key)
2532 {
2533 key->ge.opt.kill_clip_distances = vs->info.clipdist_mask & ~sctx->queued.named.rasterizer->clip_plane_enable;
2534
2535 /* Find out which VS outputs aren't used by the PS. */
2536 uint64_t outputs_written = vs->info.outputs_written_before_ps;
2537 uint64_t linked = outputs_written & sctx->ps_inputs_read_or_disabled;
2538
2539 key->ge.opt.kill_layer = vs->info.writes_layer &&
2540 sctx->framebuffer.state.layers <= 1;
2541 key->ge.opt.kill_outputs = ~linked & outputs_written;
2542 key->ge.opt.ngg_culling = sctx->ngg_culling;
2543 key->ge.mono.u.vs_export_prim_id = vs->stage != MESA_SHADER_GEOMETRY &&
2544 sctx->shader.ps.cso && sctx->shader.ps.cso->info.uses_primid;
2545
2546 if (vs->info.enabled_streamout_buffer_mask) {
2547 if (sctx->streamout.enabled_mask) {
2548 key->ge.opt.remove_streamout = 0;
2549 key->ge.opt.ngg_vs_streamout_num_verts_per_prim =
2550 sctx->gfx_level >= GFX11 ? sctx->streamout.num_verts_per_prim : 0;
2551 } else {
2552 key->ge.opt.remove_streamout = 1;
2553 key->ge.opt.ngg_vs_streamout_num_verts_per_prim = 0;
2554 }
2555 } else {
2556 key->ge.opt.remove_streamout = 0;
2557 key->ge.opt.ngg_vs_streamout_num_verts_per_prim = 0;
2558 }
2559
2560 if (sctx->gfx_level >= GFX12)
2561 key->ge.mono.remove_streamout = key->ge.opt.remove_streamout;
2562 }
2563
si_clear_vs_key_outputs(struct si_context * sctx,struct si_shader_selector * vs,union si_shader_key * key)2564 static void si_clear_vs_key_outputs(struct si_context *sctx, struct si_shader_selector *vs,
2565 union si_shader_key *key)
2566 {
2567 key->ge.opt.kill_clip_distances = 0;
2568 key->ge.opt.kill_outputs = 0;
2569 key->ge.opt.remove_streamout = 0;
2570 key->ge.opt.ngg_culling = 0;
2571 key->ge.opt.ngg_vs_streamout_num_verts_per_prim = 0;
2572 key->ge.mono.u.vs_export_prim_id = 0;
2573 key->ge.mono.remove_streamout = 0;
2574 }
2575
si_ps_key_update_framebuffer(struct si_context * sctx)2576 void si_ps_key_update_framebuffer(struct si_context *sctx)
2577 {
2578 struct si_shader_selector *sel = sctx->shader.ps.cso;
2579 union si_shader_key *key = &sctx->shader.ps.key;
2580
2581 if (!sel)
2582 return;
2583
2584 /* ps_uses_fbfetch is true only if the color buffer is bound. */
2585 if (sctx->ps_uses_fbfetch) {
2586 struct pipe_surface *cb0 = sctx->framebuffer.state.cbufs[0];
2587 struct pipe_resource *tex = cb0->texture;
2588
2589 /* 1D textures are allocated and used as 2D on GFX9. */
2590 key->ps.mono.fbfetch_msaa = sctx->framebuffer.nr_samples > 1;
2591 key->ps.mono.fbfetch_is_1D =
2592 sctx->gfx_level != GFX9 &&
2593 (tex->target == PIPE_TEXTURE_1D || tex->target == PIPE_TEXTURE_1D_ARRAY);
2594 key->ps.mono.fbfetch_layered =
2595 tex->target == PIPE_TEXTURE_1D_ARRAY || tex->target == PIPE_TEXTURE_2D_ARRAY ||
2596 tex->target == PIPE_TEXTURE_CUBE || tex->target == PIPE_TEXTURE_CUBE_ARRAY ||
2597 tex->target == PIPE_TEXTURE_3D;
2598 } else {
2599 key->ps.mono.fbfetch_msaa = 0;
2600 key->ps.mono.fbfetch_is_1D = 0;
2601 key->ps.mono.fbfetch_layered = 0;
2602 }
2603 }
2604
si_ps_key_update_framebuffer_blend_dsa_rasterizer(struct si_context * sctx)2605 void si_ps_key_update_framebuffer_blend_dsa_rasterizer(struct si_context *sctx)
2606 {
2607 struct si_shader_selector *sel = sctx->shader.ps.cso;
2608 if (!sel)
2609 return;
2610
2611 union si_shader_key *key = &sctx->shader.ps.key;
2612 struct si_state_blend *blend = sctx->queued.named.blend;
2613 struct si_state_dsa *dsa = sctx->queued.named.dsa;
2614 struct si_state_rasterizer *rs = sctx->queued.named.rasterizer;
2615 bool alpha_to_coverage = sel->info.colors_written & 0x1 && blend->alpha_to_coverage &&
2616 rs->multisample_enable && sctx->framebuffer.nr_samples >= 2;
2617 unsigned need_src_alpha_4bit = blend->need_src_alpha_4bit;
2618
2619 /* Old key data for comparison. */
2620 struct si_ps_epilog_bits old_epilog;
2621 memcpy(&old_epilog, &key->ps.part.epilog, sizeof(old_epilog));
2622 bool old_prefer_mono = key->ps.opt.prefer_mono;
2623 #ifndef NDEBUG
2624 struct si_shader_key_ps old_key;
2625 memcpy(&old_key, &key->ps, sizeof(old_key));
2626 #endif
2627
2628 key->ps.part.epilog.kill_z = sel->info.writes_z &&
2629 (!sctx->framebuffer.state.zsbuf || !dsa->depth_enabled ||
2630 (sel->info.output_z_equals_input_z && !rs->multisample_enable));
2631 key->ps.part.epilog.kill_stencil = sel->info.writes_stencil &&
2632 (!sctx->framebuffer.has_stencil || !dsa->stencil_enabled);
2633
2634 /* Remove the gl_SampleMask fragment shader output if MSAA is disabled.
2635 * This is required for correctness and it's also an optimization.
2636 */
2637 key->ps.part.epilog.kill_samplemask = sel->info.writes_samplemask &&
2638 (sctx->framebuffer.nr_samples <= 1 ||
2639 !rs->multisample_enable);
2640
2641 key->ps.part.epilog.alpha_to_one = sel->info.colors_written & 0x1 && blend->alpha_to_one &&
2642 rs->multisample_enable;
2643 /* GFX11+ always exports alpha for alpha-to-coverage via mrtz. */
2644 key->ps.part.epilog.alpha_to_coverage_via_mrtz =
2645 alpha_to_coverage && (sctx->gfx_level >= GFX11 || key->ps.part.epilog.alpha_to_one) &&
2646 ((sel->info.writes_z && !key->ps.part.epilog.kill_z) ||
2647 (sel->info.writes_stencil && !key->ps.part.epilog.kill_stencil) ||
2648 (sel->info.writes_samplemask && !key->ps.part.epilog.kill_samplemask) ||
2649 /* If both alpha-to-coverage and alpha-to-one are enabled, alpha for alpha-to-coverage must
2650 * be exported from mrtz because mrt0.a must contain 1.0 for alpha-to-one. */
2651 key->ps.part.epilog.alpha_to_one);
2652
2653 /* If alpha-to-coverage isn't exported via MRTZ, set that we need to export alpha
2654 * through MRT0.
2655 */
2656 if (alpha_to_coverage && !key->ps.part.epilog.alpha_to_coverage_via_mrtz)
2657 need_src_alpha_4bit |= 0xf;
2658
2659 /* Select the shader color format based on whether
2660 * blending or alpha are needed.
2661 */
2662 key->ps.part.epilog.spi_shader_col_format =
2663 (blend->blend_enable_4bit & need_src_alpha_4bit &
2664 sctx->framebuffer.spi_shader_col_format_blend_alpha) |
2665 (blend->blend_enable_4bit & ~need_src_alpha_4bit &
2666 sctx->framebuffer.spi_shader_col_format_blend) |
2667 (~blend->blend_enable_4bit & need_src_alpha_4bit &
2668 sctx->framebuffer.spi_shader_col_format_alpha) |
2669 (~blend->blend_enable_4bit & ~need_src_alpha_4bit &
2670 sctx->framebuffer.spi_shader_col_format);
2671 key->ps.part.epilog.spi_shader_col_format &= blend->cb_target_enabled_4bit;
2672
2673 key->ps.part.epilog.dual_src_blend_swizzle = sctx->gfx_level >= GFX11 &&
2674 blend->dual_src_blend &&
2675 (sel->info.colors_written_4bit & 0xff) == 0xff;
2676
2677 /* The output for dual source blending should have
2678 * the same format as the first output.
2679 */
2680 if (blend->dual_src_blend) {
2681 key->ps.part.epilog.spi_shader_col_format |=
2682 (key->ps.part.epilog.spi_shader_col_format & 0xf) << 4;
2683 }
2684
2685 /* If alpha-to-coverage is enabled, we have to export alpha
2686 * even if there is no color buffer.
2687 *
2688 * Gfx11 exports alpha-to-coverage via MRTZ if MRTZ is present.
2689 */
2690 if (!(key->ps.part.epilog.spi_shader_col_format & 0xf) && alpha_to_coverage &&
2691 !key->ps.part.epilog.alpha_to_coverage_via_mrtz)
2692 key->ps.part.epilog.spi_shader_col_format |= V_028710_SPI_SHADER_32_AR;
2693
2694 /* On GFX6 and GFX7 except Hawaii, the CB doesn't clamp outputs
2695 * to the range supported by the type if a channel has less
2696 * than 16 bits and the export format is 16_ABGR.
2697 */
2698 if (sctx->gfx_level <= GFX7 && sctx->family != CHIP_HAWAII) {
2699 key->ps.part.epilog.color_is_int8 = sctx->framebuffer.color_is_int8;
2700 key->ps.part.epilog.color_is_int10 = sctx->framebuffer.color_is_int10;
2701 }
2702
2703 /* Disable unwritten outputs (if WRITE_ALL_CBUFS isn't enabled). */
2704 if (!sel->info.color0_writes_all_cbufs) {
2705 key->ps.part.epilog.spi_shader_col_format &= sel->info.colors_written_4bit;
2706 key->ps.part.epilog.color_is_int8 &= sel->info.colors_written;
2707 key->ps.part.epilog.color_is_int10 &= sel->info.colors_written;
2708 }
2709
2710 /* Enable RB+ for depth-only rendering. Registers must be programmed as follows:
2711 * CB_COLOR_CONTROL.MODE = CB_DISABLE
2712 * CB_COLOR0_INFO.FORMAT = COLOR_32
2713 * CB_COLOR0_INFO.NUMBER_TYPE = NUMBER_FLOAT
2714 * SPI_SHADER_COL_FORMAT.COL0_EXPORT_FORMAT = SPI_SHADER_32_R
2715 * SX_PS_DOWNCONVERT.MRT0 = SX_RT_EXPORT_32_R
2716 *
2717 * Also, the following conditions must be met.
2718 */
2719 key->ps.part.epilog.rbplus_depth_only_opt =
2720 sctx->screen->info.rbplus_allowed &&
2721 blend->cb_target_enabled_4bit == 0 && /* implies CB_DISABLE */
2722 !alpha_to_coverage &&
2723 !sel->info.base.writes_memory &&
2724 !key->ps.part.epilog.spi_shader_col_format;
2725
2726 /* Compile PS monolithically if it eliminates code or improves performance. */
2727 if (sel->info.colors_written_4bit &
2728 /* Dual source blending never has color buffer 1 enabled, so ignore it. */
2729 (blend->dual_src_blend ? 0xffffff0f : 0xffffffff) &
2730 ~(sctx->framebuffer.colorbuf_enabled_4bit & blend->cb_target_enabled_4bit)) {
2731 /* Eliminate shader code computing the color outputs that have missing color buffer
2732 * attachments or are disabled by colormask.
2733 */
2734 key->ps.opt.prefer_mono = 1;
2735 } else if (sctx->gfx_level >= GFX11 && sel->info.base.writes_memory) {
2736 /* On gfx11, pixel shaders that write memory should be compiled with an inlined epilog,
2737 * so that the compiler can see s_endpgm and deallocates VGPRs before memory stores return.
2738 */
2739 key->ps.opt.prefer_mono = 1;
2740 } else if (key->ps.part.epilog.kill_z || key->ps.part.epilog.kill_stencil ||
2741 key->ps.part.epilog.kill_samplemask) {
2742 /* Eliminate shader code computing the Z/S/samplemask outputs. */
2743 key->ps.opt.prefer_mono = 1;
2744 } else {
2745 key->ps.opt.prefer_mono = 0;
2746 }
2747
2748 /* Update shaders only if the key changed. */
2749 if (memcmp(&key->ps.part.epilog, &old_epilog, sizeof(old_epilog)) ||
2750 key->ps.opt.prefer_mono != old_prefer_mono) {
2751 sctx->do_update_shaders = true;
2752 } else {
2753 assert(memcmp(&key->ps, &old_key, sizeof(old_key)) == 0);
2754 }
2755 }
2756
si_ps_key_update_rasterizer(struct si_context * sctx)2757 void si_ps_key_update_rasterizer(struct si_context *sctx)
2758 {
2759 struct si_shader_selector *sel = sctx->shader.ps.cso;
2760 union si_shader_key *key = &sctx->shader.ps.key;
2761 struct si_state_rasterizer *rs = sctx->queued.named.rasterizer;
2762
2763 if (!sel)
2764 return;
2765
2766 bool old_flatshade_colors = key->ps.part.prolog.flatshade_colors;
2767 bool old_clamp_color = key->ps.part.epilog.clamp_color;
2768
2769 key->ps.part.prolog.flatshade_colors = rs->flatshade && sel->info.uses_interp_color;
2770 key->ps.part.epilog.clamp_color = rs->clamp_fragment_color;
2771
2772 if (key->ps.part.prolog.flatshade_colors != old_flatshade_colors ||
2773 key->ps.part.epilog.clamp_color != old_clamp_color)
2774 sctx->do_update_shaders = true;
2775 }
2776
si_ps_key_update_dsa(struct si_context * sctx)2777 void si_ps_key_update_dsa(struct si_context *sctx)
2778 {
2779 union si_shader_key *key = &sctx->shader.ps.key;
2780
2781 key->ps.part.epilog.alpha_func = sctx->queued.named.dsa->alpha_func;
2782 }
2783
si_ps_key_update_sample_shading(struct si_context * sctx)2784 void si_ps_key_update_sample_shading(struct si_context *sctx)
2785 {
2786 struct si_shader_selector *sel = sctx->shader.ps.cso;
2787 if (!sel)
2788 return;
2789
2790 union si_shader_key *key = &sctx->shader.ps.key;
2791 unsigned ps_iter_samples = si_get_ps_iter_samples(sctx);
2792 assert(ps_iter_samples <= MAX2(1, sctx->framebuffer.nr_color_samples));
2793
2794 if (ps_iter_samples > 1 && sel->info.reads_samplemask) {
2795 /* Set samplemask_log_ps_iter=3 if full sample shading is enabled even for 2x and 4x MSAA
2796 * to get the fast path that fully replaces sample_mask_in with sample_id.
2797 */
2798 if (ps_iter_samples == sctx->framebuffer.nr_color_samples)
2799 key->ps.part.prolog.samplemask_log_ps_iter = 3;
2800 else
2801 key->ps.part.prolog.samplemask_log_ps_iter = util_logbase2(ps_iter_samples);
2802 } else {
2803 key->ps.part.prolog.samplemask_log_ps_iter = 0;
2804 }
2805 }
2806
si_ps_key_update_framebuffer_rasterizer_sample_shading(struct si_context * sctx)2807 void si_ps_key_update_framebuffer_rasterizer_sample_shading(struct si_context *sctx)
2808 {
2809 struct si_shader_selector *sel = sctx->shader.ps.cso;
2810 union si_shader_key *key = &sctx->shader.ps.key;
2811 struct si_state_rasterizer *rs = sctx->queued.named.rasterizer;
2812
2813 if (!sel)
2814 return;
2815
2816 /* Old key data for comparison. */
2817 struct si_ps_prolog_bits old_prolog;
2818 memcpy(&old_prolog, &key->ps.part.prolog, sizeof(old_prolog));
2819 bool old_interpolate_at_sample_force_center = key->ps.mono.interpolate_at_sample_force_center;
2820
2821 bool uses_persp_center = sel->info.uses_persp_center ||
2822 (!rs->flatshade && sel->info.uses_persp_center_color);
2823 bool uses_persp_centroid = sel->info.uses_persp_centroid ||
2824 (!rs->flatshade && sel->info.uses_persp_centroid_color);
2825 bool uses_persp_sample = sel->info.uses_persp_sample ||
2826 (!rs->flatshade && sel->info.uses_persp_sample_color);
2827
2828 if (!sel->info.base.fs.uses_sample_shading && rs->multisample_enable &&
2829 sctx->framebuffer.nr_samples > 1 && sctx->ps_iter_samples > 1) {
2830 key->ps.part.prolog.force_persp_sample_interp =
2831 uses_persp_center || uses_persp_centroid;
2832
2833 key->ps.part.prolog.force_linear_sample_interp =
2834 sel->info.uses_linear_center || sel->info.uses_linear_centroid;
2835
2836 key->ps.part.prolog.force_persp_center_interp = 0;
2837 key->ps.part.prolog.force_linear_center_interp = 0;
2838 key->ps.part.prolog.bc_optimize_for_persp = 0;
2839 key->ps.part.prolog.bc_optimize_for_linear = 0;
2840 key->ps.part.prolog.force_samplemask_to_helper_invocation = 0;
2841 /* Note that interpolateAt* requires center barycentrics while the PS prolog forces
2842 * per-sample barycentrics in center VGPRs, so it breaks it. The workaround is to
2843 * force monolithic compilation, which does the right thing.
2844 */
2845 key->ps.mono.force_mono = sel->info.uses_interp_at_offset || sel->info.uses_interp_at_sample;
2846 key->ps.mono.interpolate_at_sample_force_center = 0;
2847 } else if (rs->multisample_enable && sctx->framebuffer.nr_samples > 1) {
2848 /* Note that sample shading is possible here. If it's enabled, all barycentrics are
2849 * already set to "sample" except at_offset/at_sample.
2850 */
2851 key->ps.part.prolog.force_persp_sample_interp = 0;
2852 key->ps.part.prolog.force_linear_sample_interp = 0;
2853 key->ps.part.prolog.force_persp_center_interp = 0;
2854 key->ps.part.prolog.force_linear_center_interp = 0;
2855 key->ps.part.prolog.bc_optimize_for_persp =
2856 uses_persp_center && uses_persp_centroid;
2857 key->ps.part.prolog.bc_optimize_for_linear =
2858 sel->info.uses_linear_center && sel->info.uses_linear_centroid;
2859 key->ps.part.prolog.get_frag_coord_from_pixel_coord =
2860 !sel->info.base.fs.uses_sample_shading && sel->info.reads_frag_coord_mask & 0x3;
2861 key->ps.part.prolog.force_samplemask_to_helper_invocation = 0;
2862 key->ps.mono.force_mono = 0;
2863 key->ps.mono.interpolate_at_sample_force_center = 0;
2864 } else {
2865 key->ps.part.prolog.force_persp_sample_interp = 0;
2866 key->ps.part.prolog.force_linear_sample_interp = 0;
2867
2868 /* Make sure SPI doesn't compute more than 1 pair
2869 * of (i,j), which is the optimization here. */
2870 key->ps.part.prolog.force_persp_center_interp = uses_persp_center +
2871 uses_persp_centroid +
2872 uses_persp_sample > 1;
2873
2874 key->ps.part.prolog.force_linear_center_interp = sel->info.uses_linear_center +
2875 sel->info.uses_linear_centroid +
2876 sel->info.uses_linear_sample > 1;
2877 key->ps.part.prolog.bc_optimize_for_persp = 0;
2878 key->ps.part.prolog.bc_optimize_for_linear = 0;
2879 key->ps.part.prolog.get_frag_coord_from_pixel_coord =
2880 !!(sel->info.reads_frag_coord_mask & 0x3);
2881 key->ps.part.prolog.force_samplemask_to_helper_invocation = sel->info.reads_samplemask;
2882 key->ps.mono.force_mono = 0;
2883 key->ps.mono.interpolate_at_sample_force_center = sel->info.uses_interp_at_sample;
2884 }
2885
2886 /* Update shaders only if the key changed. */
2887 if (memcmp(&key->ps.part.prolog, &old_prolog, sizeof(old_prolog)) ||
2888 key->ps.mono.interpolate_at_sample_force_center != old_interpolate_at_sample_force_center)
2889 sctx->do_update_shaders = true;
2890 }
2891
2892 /* Compute the key for the hw shader variant */
si_shader_selector_key(struct pipe_context * ctx,struct si_shader_selector * sel,union si_shader_key * key)2893 static inline void si_shader_selector_key(struct pipe_context *ctx, struct si_shader_selector *sel,
2894 union si_shader_key *key)
2895 {
2896 struct si_context *sctx = (struct si_context *)ctx;
2897
2898 switch (sel->stage) {
2899 case MESA_SHADER_VERTEX:
2900 if (!sctx->shader.tes.cso && !sctx->shader.gs.cso)
2901 si_get_vs_key_outputs(sctx, sel, key);
2902 else
2903 si_clear_vs_key_outputs(sctx, sel, key);
2904 break;
2905 case MESA_SHADER_TESS_CTRL:
2906 if (sctx->gfx_level >= GFX9) {
2907 si_get_vs_key_inputs(sctx, key);
2908 key->ge.part.tcs.ls = sctx->shader.vs.cso;
2909 }
2910 break;
2911 case MESA_SHADER_TESS_EVAL:
2912 if (!sctx->shader.gs.cso)
2913 si_get_vs_key_outputs(sctx, sel, key);
2914 else
2915 si_clear_vs_key_outputs(sctx, sel, key);
2916 break;
2917 case MESA_SHADER_GEOMETRY:
2918 if (sctx->gfx_level >= GFX9) {
2919 if (sctx->shader.tes.cso) {
2920 si_clear_vs_key_inputs(key);
2921 key->ge.part.gs.es = sctx->shader.tes.cso;
2922 } else {
2923 si_get_vs_key_inputs(sctx, key);
2924 key->ge.part.gs.es = sctx->shader.vs.cso;
2925 }
2926
2927 /* Only NGG can eliminate GS outputs, because the code is shared with VS. */
2928 if (sctx->ngg)
2929 si_get_vs_key_outputs(sctx, sel, key);
2930 else
2931 si_clear_vs_key_outputs(sctx, sel, key);
2932 }
2933 break;
2934 case MESA_SHADER_FRAGMENT:
2935 break;
2936 default:
2937 assert(0);
2938 }
2939 }
2940
si_build_shader_variant(struct si_shader * shader,int thread_index,bool low_priority)2941 static void si_build_shader_variant(struct si_shader *shader, int thread_index, bool low_priority)
2942 {
2943 struct si_shader_selector *sel = shader->selector;
2944 struct si_screen *sscreen = sel->screen;
2945 struct ac_llvm_compiler **compiler;
2946 struct util_debug_callback *debug = &shader->compiler_ctx_state.debug;
2947
2948 if (thread_index >= 0) {
2949 if (low_priority) {
2950 assert(thread_index < (int)ARRAY_SIZE(sscreen->compiler_lowp));
2951 compiler = &sscreen->compiler_lowp[thread_index];
2952 } else {
2953 assert(thread_index < (int)ARRAY_SIZE(sscreen->compiler));
2954 compiler = &sscreen->compiler[thread_index];
2955 }
2956 if (!debug->async)
2957 debug = NULL;
2958 } else {
2959 assert(!low_priority);
2960 compiler = &shader->compiler_ctx_state.compiler;
2961 }
2962
2963 if (!sel->info.base.use_aco_amd && !*compiler)
2964 *compiler = si_create_llvm_compiler(sscreen);
2965
2966 if (unlikely(!si_create_shader_variant(sscreen, *compiler, shader, debug))) {
2967 PRINT_ERR("Failed to build shader variant (type=%u)\n", sel->stage);
2968 shader->compilation_failed = true;
2969 return;
2970 }
2971
2972 if (shader->compiler_ctx_state.is_debug_context) {
2973 FILE *f = open_memstream(&shader->shader_log, &shader->shader_log_size);
2974 if (f) {
2975 si_shader_dump(sscreen, shader, NULL, f, false);
2976 fclose(f);
2977 }
2978 }
2979
2980 si_shader_init_pm4_state(sscreen, shader);
2981 }
2982
si_build_shader_variant_low_priority(void * job,void * gdata,int thread_index)2983 static void si_build_shader_variant_low_priority(void *job, void *gdata, int thread_index)
2984 {
2985 struct si_shader *shader = (struct si_shader *)job;
2986
2987 assert(thread_index >= 0);
2988
2989 si_build_shader_variant(shader, thread_index, true);
2990 }
2991
2992 /* This should be const, but C++ doesn't allow implicit zero-initialization with const. */
2993 static union si_shader_key zeroed;
2994
si_check_missing_main_part(struct si_screen * sscreen,struct si_shader_selector * sel,struct si_compiler_ctx_state * compiler_state,const union si_shader_key * key,unsigned wave_size)2995 static bool si_check_missing_main_part(struct si_screen *sscreen, struct si_shader_selector *sel,
2996 struct si_compiler_ctx_state *compiler_state,
2997 const union si_shader_key *key, unsigned wave_size)
2998 {
2999 struct si_shader **mainp = si_get_main_shader_part(sel, key, wave_size);
3000
3001 if (!*mainp) {
3002 struct si_shader *main_part = CALLOC_STRUCT(si_shader);
3003
3004 if (!main_part)
3005 return false;
3006
3007 /* We can leave the fence as permanently signaled because the
3008 * main part becomes visible globally only after it has been
3009 * compiled. */
3010 util_queue_fence_init(&main_part->ready);
3011
3012 main_part->selector = sel;
3013 if (sel->stage <= MESA_SHADER_GEOMETRY) {
3014 main_part->key.ge.as_es = key->ge.as_es;
3015 main_part->key.ge.as_ls = key->ge.as_ls;
3016 main_part->key.ge.as_ngg = key->ge.as_ngg;
3017 }
3018 main_part->is_monolithic = false;
3019 main_part->wave_size = wave_size;
3020
3021 if (!si_compile_shader(sscreen, compiler_state->compiler, main_part,
3022 &compiler_state->debug)) {
3023 FREE(main_part);
3024 return false;
3025 }
3026 *mainp = main_part;
3027 }
3028 return true;
3029 }
3030
3031 /* A helper to copy *key to *local_key and return local_key. */
3032 template<typename SHADER_KEY_TYPE>
3033 static ALWAYS_INLINE const SHADER_KEY_TYPE *
use_local_key_copy(const SHADER_KEY_TYPE * key,SHADER_KEY_TYPE * local_key,unsigned key_size)3034 use_local_key_copy(const SHADER_KEY_TYPE *key, SHADER_KEY_TYPE *local_key, unsigned key_size)
3035 {
3036 if (key != local_key)
3037 memcpy(local_key, key, key_size);
3038
3039 return local_key;
3040 }
3041
3042 #define NO_INLINE_UNIFORMS false
3043
3044 /**
3045 * Select a shader variant according to the shader key.
3046 *
3047 * This uses a C++ template to compute the optimal memcmp size at compile time, which is important
3048 * for getting inlined memcmp. The memcmp size depends on the shader key type and whether inlined
3049 * uniforms are enabled.
3050 */
3051 template<bool INLINE_UNIFORMS = true, typename SHADER_KEY_TYPE>
si_shader_select_with_key(struct si_context * sctx,struct si_shader_ctx_state * state,const SHADER_KEY_TYPE * key)3052 static int si_shader_select_with_key(struct si_context *sctx, struct si_shader_ctx_state *state,
3053 const SHADER_KEY_TYPE *key)
3054 {
3055 struct si_screen *sscreen = sctx->screen;
3056 struct si_shader_selector *sel = state->cso;
3057 struct si_shader_selector *previous_stage_sel = NULL;
3058 struct si_shader *current = state->current;
3059 struct si_shader *shader = NULL;
3060 const SHADER_KEY_TYPE *zeroed_key = (SHADER_KEY_TYPE*)&zeroed;
3061
3062 /* "opt" must be the last field and "inlined_uniform_values" must be the last field inside opt.
3063 * If there is padding, insert the padding manually before opt or inside opt.
3064 */
3065 STATIC_ASSERT(offsetof(SHADER_KEY_TYPE, opt) + sizeof(key->opt) == sizeof(*key));
3066 STATIC_ASSERT(offsetof(SHADER_KEY_TYPE, opt.inlined_uniform_values) +
3067 sizeof(key->opt.inlined_uniform_values) == sizeof(*key));
3068
3069 const unsigned key_size_no_uniforms = sizeof(*key) - sizeof(key->opt.inlined_uniform_values);
3070 /* Don't compare inlined_uniform_values if uniform inlining is disabled. */
3071 const unsigned key_size = INLINE_UNIFORMS ? sizeof(*key) : key_size_no_uniforms;
3072 const unsigned key_opt_size =
3073 INLINE_UNIFORMS ? sizeof(key->opt) :
3074 sizeof(key->opt) - sizeof(key->opt.inlined_uniform_values);
3075
3076 /* si_shader_select_with_key must not modify 'key' because it would affect future shaders.
3077 * If we need to modify it for this specific shader (eg: to disable optimizations), we
3078 * use a copy.
3079 */
3080 SHADER_KEY_TYPE local_key;
3081
3082 if (unlikely(sscreen->debug_flags & DBG(NO_OPT_VARIANT))) {
3083 /* Disable shader variant optimizations. */
3084 key = use_local_key_copy<SHADER_KEY_TYPE>(key, &local_key, key_size);
3085 memset(&local_key.opt, 0, key_opt_size);
3086 }
3087
3088 again:
3089 /* Check if we don't need to change anything.
3090 * This path is also used for most shaders that don't need multiple
3091 * variants, it will cost just a computation of the key and this
3092 * test. */
3093 if (likely(current && memcmp(¤t->key, key, key_size) == 0)) {
3094 if (unlikely(!util_queue_fence_is_signalled(¤t->ready))) {
3095 if (current->is_optimized) {
3096 key = use_local_key_copy(key, &local_key, key_size);
3097 memset(&local_key.opt, 0, key_opt_size);
3098 goto current_not_ready;
3099 }
3100
3101 util_queue_fence_wait(¤t->ready);
3102 }
3103
3104 return current->compilation_failed ? -1 : 0;
3105 }
3106 current_not_ready:
3107
3108 /* This must be done before the mutex is locked, because async GS
3109 * compilation calls this function too, and therefore must enter
3110 * the mutex first.
3111 */
3112 util_queue_fence_wait(&sel->ready);
3113
3114 simple_mtx_lock(&sel->mutex);
3115
3116 int variant_count = 0;
3117 const int max_inline_uniforms_variants = 5;
3118
3119 /* Find the shader variant. */
3120 const unsigned cnt = sel->variants_count;
3121 for (unsigned i = 0; i < cnt; i++) {
3122 const SHADER_KEY_TYPE *iter_key = (const SHADER_KEY_TYPE *)&sel->keys[i];
3123
3124 if (memcmp(iter_key, key, key_size_no_uniforms) == 0) {
3125 struct si_shader *iter = sel->variants[i];
3126
3127 /* Check the inlined uniform values separately, and count
3128 * the number of variants based on them.
3129 */
3130 if (key->opt.inline_uniforms &&
3131 memcmp(iter_key->opt.inlined_uniform_values,
3132 key->opt.inlined_uniform_values,
3133 MAX_INLINABLE_UNIFORMS * 4) != 0) {
3134 if (variant_count++ > max_inline_uniforms_variants) {
3135 key = use_local_key_copy(key, &local_key, key_size);
3136 /* Too many variants. Disable inlining for this shader. */
3137 local_key.opt.inline_uniforms = 0;
3138 memset(local_key.opt.inlined_uniform_values, 0, MAX_INLINABLE_UNIFORMS * 4);
3139 simple_mtx_unlock(&sel->mutex);
3140 goto again;
3141 }
3142 continue;
3143 }
3144
3145 simple_mtx_unlock(&sel->mutex);
3146
3147 if (unlikely(!util_queue_fence_is_signalled(&iter->ready))) {
3148 /* If it's an optimized shader and its compilation has
3149 * been started but isn't done, use the unoptimized
3150 * shader so as not to cause a stall due to compilation.
3151 */
3152 if (iter->is_optimized) {
3153 key = use_local_key_copy(key, &local_key, key_size);
3154 memset(&local_key.opt, 0, key_opt_size);
3155 goto again;
3156 }
3157
3158 util_queue_fence_wait(&iter->ready);
3159 }
3160
3161 if (iter->compilation_failed) {
3162 return -1; /* skip the draw call */
3163 }
3164
3165 state->current = sel->variants[i];
3166 return 0;
3167 }
3168 }
3169
3170 /* Build a new shader. */
3171 shader = CALLOC_STRUCT(si_shader);
3172 if (!shader) {
3173 simple_mtx_unlock(&sel->mutex);
3174 return -ENOMEM;
3175 }
3176
3177 util_queue_fence_init(&shader->ready);
3178
3179 if (!sel->info.base.use_aco_amd && !sctx->compiler)
3180 sctx->compiler = si_create_llvm_compiler(sctx->screen);
3181
3182 shader->selector = sel;
3183 *((SHADER_KEY_TYPE*)&shader->key) = *key;
3184 shader->wave_size = si_determine_wave_size(sscreen, shader);
3185 shader->compiler_ctx_state.compiler = sctx->compiler;
3186 shader->compiler_ctx_state.debug = sctx->debug;
3187 shader->compiler_ctx_state.is_debug_context = sctx->is_debug;
3188
3189 /* If this is a merged shader, get the first shader's selector. */
3190 if (sscreen->info.gfx_level >= GFX9) {
3191 if (sel->stage == MESA_SHADER_TESS_CTRL)
3192 previous_stage_sel = ((struct si_shader_key_ge*)key)->part.tcs.ls;
3193 else if (sel->stage == MESA_SHADER_GEOMETRY)
3194 previous_stage_sel = ((struct si_shader_key_ge*)key)->part.gs.es;
3195
3196 /* We need to wait for the previous shader. */
3197 if (previous_stage_sel)
3198 util_queue_fence_wait(&previous_stage_sel->ready);
3199 }
3200
3201 bool is_pure_monolithic =
3202 sscreen->use_monolithic_shaders || memcmp(&key->mono, &zeroed_key->mono, sizeof(key->mono)) != 0;
3203
3204 /* Compile the main shader part if it doesn't exist. This can happen
3205 * if the initial guess was wrong.
3206 */
3207 if (!is_pure_monolithic) {
3208 bool ok = true;
3209
3210 /* Make sure the main shader part is present. This is needed
3211 * for shaders that can be compiled as VS, LS, or ES, and only
3212 * one of them is compiled at creation.
3213 *
3214 * It is also needed for GS, which can be compiled as non-NGG
3215 * and NGG.
3216 *
3217 * For merged shaders, check that the starting shader's main
3218 * part is present.
3219 */
3220 if (previous_stage_sel) {
3221 union si_shader_key shader1_key = zeroed;
3222
3223 if (sel->stage == MESA_SHADER_TESS_CTRL) {
3224 shader1_key.ge.as_ls = 1;
3225 } else if (sel->stage == MESA_SHADER_GEOMETRY) {
3226 shader1_key.ge.as_es = 1;
3227 shader1_key.ge.as_ngg = ((struct si_shader_key_ge*)key)->as_ngg; /* for Wave32 vs Wave64 */
3228 } else {
3229 assert(0);
3230 }
3231
3232 simple_mtx_lock(&previous_stage_sel->mutex);
3233 ok = si_check_missing_main_part(sscreen, previous_stage_sel, &shader->compiler_ctx_state,
3234 &shader1_key, shader->wave_size);
3235 simple_mtx_unlock(&previous_stage_sel->mutex);
3236 }
3237
3238 if (ok) {
3239 ok = si_check_missing_main_part(sscreen, sel, &shader->compiler_ctx_state,
3240 (union si_shader_key*)key, shader->wave_size);
3241 }
3242
3243 if (!ok) {
3244 FREE(shader);
3245 simple_mtx_unlock(&sel->mutex);
3246 return -ENOMEM; /* skip the draw call */
3247 }
3248 }
3249
3250 if (sel->variants_count == sel->variants_max_count) {
3251 sel->variants_max_count += 2;
3252 sel->variants = (struct si_shader**)
3253 realloc(sel->variants, sel->variants_max_count * sizeof(struct si_shader*));
3254 sel->keys = (union si_shader_key*)
3255 realloc(sel->keys, sel->variants_max_count * sizeof(union si_shader_key));
3256 }
3257
3258 /* Keep the reference to the 1st shader of merged shaders, so that
3259 * Gallium can't destroy it before we destroy the 2nd shader.
3260 *
3261 * Set sctx = NULL, because it's unused if we're not releasing
3262 * the shader, and we don't have any sctx here.
3263 */
3264 si_shader_selector_reference(NULL, &shader->previous_stage_sel, previous_stage_sel);
3265
3266 /* Monolithic-only shaders don't make a distinction between optimized
3267 * and unoptimized. */
3268 shader->is_monolithic =
3269 is_pure_monolithic || memcmp(&key->opt, &zeroed_key->opt, key_opt_size) != 0;
3270
3271 shader->is_optimized = !is_pure_monolithic &&
3272 memcmp(&key->opt, &zeroed_key->opt, key_opt_size) != 0;
3273
3274 /* If it's an optimized shader, compile it asynchronously. */
3275 if (shader->is_optimized) {
3276 /* Compile it asynchronously. */
3277 util_queue_add_job(&sscreen->shader_compiler_queue_opt_variants, shader, &shader->ready,
3278 si_build_shader_variant_low_priority, NULL, 0);
3279
3280 /* Add only after the ready fence was reset, to guard against a
3281 * race with si_bind_XX_shader. */
3282 sel->variants[sel->variants_count] = shader;
3283 sel->keys[sel->variants_count] = shader->key;
3284 sel->variants_count++;
3285
3286 /* Use the default (unoptimized) shader for now. */
3287 key = use_local_key_copy(key, &local_key, key_size);
3288 memset(&local_key.opt, 0, key_opt_size);
3289 simple_mtx_unlock(&sel->mutex);
3290
3291 if (sscreen->options.sync_compile)
3292 util_queue_fence_wait(&shader->ready);
3293
3294 goto again;
3295 }
3296
3297 /* Reset the fence before adding to the variant list. */
3298 util_queue_fence_reset(&shader->ready);
3299
3300 sel->variants[sel->variants_count] = shader;
3301 sel->keys[sel->variants_count] = shader->key;
3302 sel->variants_count++;
3303
3304 simple_mtx_unlock(&sel->mutex);
3305
3306 assert(!shader->is_optimized);
3307 si_build_shader_variant(shader, -1, false);
3308
3309 util_queue_fence_signal(&shader->ready);
3310
3311 if (!shader->compilation_failed)
3312 state->current = shader;
3313
3314 return shader->compilation_failed ? -1 : 0;
3315 }
3316
si_shader_select(struct pipe_context * ctx,struct si_shader_ctx_state * state)3317 int si_shader_select(struct pipe_context *ctx, struct si_shader_ctx_state *state)
3318 {
3319 struct si_context *sctx = (struct si_context *)ctx;
3320
3321 si_shader_selector_key(ctx, state->cso, &state->key);
3322
3323 if (state->cso->stage == MESA_SHADER_FRAGMENT) {
3324 if (state->key.ps.opt.inline_uniforms)
3325 return si_shader_select_with_key(sctx, state, &state->key.ps);
3326 else
3327 return si_shader_select_with_key<NO_INLINE_UNIFORMS>(sctx, state, &state->key.ps);
3328 } else {
3329 if (state->key.ge.opt.inline_uniforms) {
3330 return si_shader_select_with_key(sctx, state, &state->key.ge);
3331 } else {
3332 return si_shader_select_with_key<NO_INLINE_UNIFORMS>(sctx, state, &state->key.ge);
3333 }
3334 }
3335 }
3336
si_parse_next_shader_property(const struct si_shader_info * info,union si_shader_key * key)3337 static void si_parse_next_shader_property(const struct si_shader_info *info,
3338 union si_shader_key *key)
3339 {
3340 gl_shader_stage next_shader = info->base.next_stage;
3341
3342 switch (info->base.stage) {
3343 case MESA_SHADER_VERTEX:
3344 switch (next_shader) {
3345 case MESA_SHADER_GEOMETRY:
3346 key->ge.as_es = 1;
3347 break;
3348 case MESA_SHADER_TESS_CTRL:
3349 case MESA_SHADER_TESS_EVAL:
3350 key->ge.as_ls = 1;
3351 break;
3352 default:
3353 /* If POSITION isn't written, it can only be a HW VS
3354 * if streamout is used. If streamout isn't used,
3355 * assume that it's a HW LS. (the next shader is TCS)
3356 * This heuristic is needed for separate shader objects.
3357 */
3358 if (!info->writes_position && !info->enabled_streamout_buffer_mask)
3359 key->ge.as_ls = 1;
3360 }
3361 break;
3362
3363 case MESA_SHADER_TESS_EVAL:
3364 if (next_shader == MESA_SHADER_GEOMETRY || !info->writes_position)
3365 key->ge.as_es = 1;
3366 break;
3367
3368 default:;
3369 }
3370 }
3371
3372 /**
3373 * Compile the main shader part or the monolithic shader as part of
3374 * si_shader_selector initialization. Since it can be done asynchronously,
3375 * there is no way to report compile failures to applications.
3376 */
si_init_shader_selector_async(void * job,void * gdata,int thread_index)3377 static void si_init_shader_selector_async(void *job, void *gdata, int thread_index)
3378 {
3379 struct si_shader_selector *sel = (struct si_shader_selector *)job;
3380 struct si_screen *sscreen = sel->screen;
3381 struct ac_llvm_compiler **compiler;
3382 struct util_debug_callback *debug = &sel->compiler_ctx_state.debug;
3383
3384 assert(!debug->debug_message || debug->async);
3385 assert(thread_index >= 0);
3386 assert(thread_index < (int)ARRAY_SIZE(sscreen->compiler));
3387 compiler = &sscreen->compiler[thread_index];
3388
3389 if (!sel->info.base.use_aco_amd && !*compiler)
3390 *compiler = si_create_llvm_compiler(sscreen);
3391
3392 /* Serialize NIR to save memory. Monolithic shader variants
3393 * have to deserialize NIR before compilation.
3394 */
3395 if (sel->nir) {
3396 struct blob blob;
3397 size_t size;
3398
3399 blob_init(&blob);
3400 /* true = remove optional debugging data to increase
3401 * the likehood of getting more shader cache hits.
3402 * It also drops variable names, so we'll save more memory.
3403 * If NIR debug prints are used we don't strip to get more
3404 * useful logs.
3405 */
3406 nir_serialize(&blob, sel->nir, NIR_DEBUG(PRINT) == 0);
3407 blob_finish_get_buffer(&blob, &sel->nir_binary, &size);
3408 sel->nir_size = size;
3409 }
3410
3411 /* Compile the main shader part for use with a prolog and/or epilog.
3412 * If this fails, the driver will try to compile a monolithic shader
3413 * on demand.
3414 */
3415 if (!sscreen->use_monolithic_shaders) {
3416 struct si_shader *shader = CALLOC_STRUCT(si_shader);
3417 unsigned char ir_sha1_cache_key[20];
3418
3419 if (!shader) {
3420 fprintf(stderr, "radeonsi: can't allocate a main shader part\n");
3421 return;
3422 }
3423
3424 /* We can leave the fence signaled because use of the default
3425 * main part is guarded by the selector's ready fence. */
3426 util_queue_fence_init(&shader->ready);
3427
3428 shader->selector = sel;
3429 shader->is_monolithic = false;
3430 si_parse_next_shader_property(&sel->info, &shader->key);
3431
3432 if (sel->stage <= MESA_SHADER_GEOMETRY &&
3433 sscreen->use_ngg && (!sel->info.enabled_streamout_buffer_mask ||
3434 sscreen->info.gfx_level >= GFX11) &&
3435 ((sel->stage == MESA_SHADER_VERTEX && !shader->key.ge.as_ls) ||
3436 sel->stage == MESA_SHADER_TESS_EVAL || sel->stage == MESA_SHADER_GEOMETRY))
3437 shader->key.ge.as_ngg = 1;
3438
3439 shader->wave_size = si_determine_wave_size(sscreen, shader);
3440
3441 if (sel->nir) {
3442 if (sel->stage <= MESA_SHADER_GEOMETRY) {
3443 si_get_ir_cache_key(sel, shader->key.ge.as_ngg, shader->key.ge.as_es,
3444 shader->wave_size, ir_sha1_cache_key);
3445 } else {
3446 si_get_ir_cache_key(sel, false, false, shader->wave_size, ir_sha1_cache_key);
3447 }
3448 }
3449
3450 /* Try to load the shader from the shader cache. */
3451 simple_mtx_lock(&sscreen->shader_cache_mutex);
3452
3453 if (si_shader_cache_load_shader(sscreen, ir_sha1_cache_key, shader)) {
3454 simple_mtx_unlock(&sscreen->shader_cache_mutex);
3455 si_shader_dump_stats_for_shader_db(sscreen, shader, debug);
3456 } else {
3457 simple_mtx_unlock(&sscreen->shader_cache_mutex);
3458
3459 /* Compile the shader if it hasn't been loaded from the cache. */
3460 if (!si_compile_shader(sscreen, *compiler, shader, debug)) {
3461 fprintf(stderr,
3462 "radeonsi: can't compile a main shader part (type: %s, name: %s).\n"
3463 "This is probably a driver bug, please report "
3464 "it to https://gitlab.freedesktop.org/mesa/mesa/-/issues.\n",
3465 gl_shader_stage_name(shader->selector->stage),
3466 shader->selector->info.base.name);
3467 FREE(shader);
3468 return;
3469 }
3470
3471 simple_mtx_lock(&sscreen->shader_cache_mutex);
3472 si_shader_cache_insert_shader(sscreen, ir_sha1_cache_key, shader, true);
3473 simple_mtx_unlock(&sscreen->shader_cache_mutex);
3474 }
3475
3476 *si_get_main_shader_part(sel, &shader->key, shader->wave_size) = shader;
3477
3478 /* Unset "outputs_written" flags for outputs converted to
3479 * DEFAULT_VAL, so that later inter-shader optimizations don't
3480 * try to eliminate outputs that don't exist in the final
3481 * shader.
3482 *
3483 * This is only done if non-monolithic shaders are enabled.
3484 */
3485 if ((sel->stage == MESA_SHADER_VERTEX ||
3486 sel->stage == MESA_SHADER_TESS_EVAL ||
3487 sel->stage == MESA_SHADER_GEOMETRY) &&
3488 !shader->key.ge.as_ls && !shader->key.ge.as_es) {
3489 unsigned i;
3490
3491 for (i = 0; i < sel->info.num_outputs; i++) {
3492 unsigned semantic = sel->info.output_semantic[i];
3493 unsigned ps_input_cntl = shader->info.vs_output_ps_input_cntl[semantic];
3494
3495 /* OFFSET=0x20 means DEFAULT_VAL, which means VS doesn't export it. */
3496 if (G_028644_OFFSET(ps_input_cntl) != 0x20)
3497 continue;
3498
3499 unsigned id;
3500
3501 /* Remove the output from the mask. */
3502 if ((semantic <= VARYING_SLOT_VAR31 || semantic >= VARYING_SLOT_VAR0_16BIT) &&
3503 semantic != VARYING_SLOT_POS &&
3504 semantic != VARYING_SLOT_PSIZ &&
3505 semantic != VARYING_SLOT_CLIP_VERTEX &&
3506 semantic != VARYING_SLOT_EDGE &&
3507 semantic != VARYING_SLOT_LAYER) {
3508 id = si_shader_io_get_unique_index(semantic);
3509 sel->info.outputs_written_before_ps &= ~(1ull << id);
3510 }
3511 }
3512 }
3513 }
3514
3515 /* Free NIR. We only keep serialized NIR after this point. */
3516 if (sel->nir) {
3517 ralloc_free(sel->nir);
3518 sel->nir = NULL;
3519 }
3520 }
3521
si_schedule_initial_compile(struct si_context * sctx,gl_shader_stage stage,struct util_queue_fence * ready_fence,struct si_compiler_ctx_state * compiler_ctx_state,void * job,util_queue_execute_func execute)3522 void si_schedule_initial_compile(struct si_context *sctx, gl_shader_stage stage,
3523 struct util_queue_fence *ready_fence,
3524 struct si_compiler_ctx_state *compiler_ctx_state, void *job,
3525 util_queue_execute_func execute)
3526 {
3527 util_queue_fence_init(ready_fence);
3528
3529 struct util_async_debug_callback async_debug;
3530 bool debug = (sctx->debug.debug_message && !sctx->debug.async) || sctx->is_debug ||
3531 si_can_dump_shader(sctx->screen, stage, SI_DUMP_ALWAYS);
3532
3533 if (debug) {
3534 u_async_debug_init(&async_debug);
3535 compiler_ctx_state->debug = async_debug.base;
3536 }
3537
3538 util_queue_add_job(&sctx->screen->shader_compiler_queue, job, ready_fence, execute, NULL, 0);
3539
3540 if (debug) {
3541 util_queue_fence_wait(ready_fence);
3542 u_async_debug_drain(&async_debug, &sctx->debug);
3543 u_async_debug_cleanup(&async_debug);
3544 }
3545
3546 if (sctx->screen->options.sync_compile)
3547 util_queue_fence_wait(ready_fence);
3548 }
3549
3550 /* Return descriptor slot usage masks from the given shader info. */
si_get_active_slot_masks(struct si_screen * sscreen,const struct si_shader_info * info,uint64_t * const_and_shader_buffers,uint64_t * samplers_and_images)3551 void si_get_active_slot_masks(struct si_screen *sscreen, const struct si_shader_info *info,
3552 uint64_t *const_and_shader_buffers, uint64_t *samplers_and_images)
3553 {
3554 unsigned start, num_shaderbufs, num_constbufs, num_images, num_msaa_images, num_samplers;
3555
3556 num_shaderbufs = info->base.num_ssbos;
3557 num_constbufs = info->base.num_ubos;
3558 /* two 8-byte images share one 16-byte slot */
3559 num_images = align(info->base.num_images, 2);
3560 num_msaa_images = align(BITSET_LAST_BIT(info->base.msaa_images), 2);
3561 num_samplers = BITSET_LAST_BIT(info->base.textures_used);
3562
3563 /* The layout is: sb[last] ... sb[0], cb[0] ... cb[last] */
3564 start = si_get_shaderbuf_slot(num_shaderbufs - 1);
3565 *const_and_shader_buffers = u_bit_consecutive64(start, num_shaderbufs + num_constbufs);
3566
3567 /* The layout is:
3568 * - fmask[last] ... fmask[0] go to [15-last .. 15]
3569 * - image[last] ... image[0] go to [31-last .. 31]
3570 * - sampler[0] ... sampler[last] go to [32 .. 32+last*2]
3571 *
3572 * FMASKs for images are placed separately, because MSAA images are rare,
3573 * and so we can benefit from a better cache hit rate if we keep image
3574 * descriptors together.
3575 */
3576 if (sscreen->info.gfx_level < GFX11 && num_msaa_images)
3577 num_images = SI_NUM_IMAGES + num_msaa_images; /* add FMASK descriptors */
3578
3579 start = si_get_image_slot(num_images - 1) / 2;
3580 *samplers_and_images = u_bit_consecutive64(start, num_images / 2 + num_samplers);
3581 }
3582
si_create_shader_selector(struct pipe_context * ctx,const struct pipe_shader_state * state)3583 static void *si_create_shader_selector(struct pipe_context *ctx,
3584 const struct pipe_shader_state *state)
3585 {
3586 struct si_screen *sscreen = (struct si_screen *)ctx->screen;
3587 struct si_context *sctx = (struct si_context *)ctx;
3588 struct si_shader_selector *sel = CALLOC_STRUCT(si_shader_selector);
3589
3590 if (!sel)
3591 return NULL;
3592
3593 sel->screen = sscreen;
3594 sel->compiler_ctx_state.debug = sctx->debug;
3595 sel->compiler_ctx_state.is_debug_context = sctx->is_debug;
3596 sel->variants_max_count = 2;
3597 sel->keys = (union si_shader_key *)
3598 realloc(NULL, sel->variants_max_count * sizeof(union si_shader_key));
3599 sel->variants = (struct si_shader **)
3600 realloc(NULL, sel->variants_max_count * sizeof(struct si_shader *));
3601
3602 if (state->type == PIPE_SHADER_IR_TGSI) {
3603 sel->nir = tgsi_to_nir(state->tokens, ctx->screen, true);
3604 } else {
3605 assert(state->type == PIPE_SHADER_IR_NIR);
3606 sel->nir = (nir_shader*)state->ir.nir;
3607 }
3608
3609 si_nir_scan_shader(sscreen, sel->nir, &sel->info, false);
3610
3611 sel->stage = sel->nir->info.stage;
3612 sel->const_and_shader_buf_descriptors_index =
3613 si_const_and_shader_buffer_descriptors_idx(sel->stage);
3614 sel->sampler_and_images_descriptors_index =
3615 si_sampler_and_image_descriptors_idx(sel->stage);
3616
3617 if (si_can_dump_shader(sscreen, sel->stage, SI_DUMP_INIT_NIR))
3618 nir_print_shader(sel->nir, stderr);
3619
3620 p_atomic_inc(&sscreen->num_shaders_created);
3621 si_get_active_slot_masks(sscreen, &sel->info, &sel->active_const_and_shader_buffers,
3622 &sel->active_samplers_and_images);
3623
3624 switch (sel->stage) {
3625 case MESA_SHADER_GEOMETRY:
3626 /* Only possibilities: POINTS, LINE_STRIP, TRIANGLES */
3627 sel->rast_prim = (enum mesa_prim)sel->nir->info.gs.output_primitive;
3628 if (util_rast_prim_is_triangles(sel->rast_prim))
3629 sel->rast_prim = MESA_PRIM_TRIANGLES;
3630
3631 /* EN_MAX_VERT_OUT_PER_GS_INSTANCE does not work with tessellation so
3632 * we can't split workgroups. Disable ngg if any of the following conditions is true:
3633 * - num_invocations * gs.vertices_out > 256
3634 * - LDS usage is too high
3635 */
3636 sel->tess_turns_off_ngg = sscreen->info.gfx_level >= GFX10 &&
3637 sscreen->info.gfx_level <= GFX10_3 &&
3638 (sel->nir->info.gs.invocations * sel->nir->info.gs.vertices_out > 256 ||
3639 sel->nir->info.gs.invocations * sel->nir->info.gs.vertices_out *
3640 (sel->info.num_outputs * 4 + 1) > 6500 /* max dw per GS primitive */);
3641 break;
3642
3643 case MESA_SHADER_VERTEX:
3644 case MESA_SHADER_TESS_EVAL:
3645 if (sel->stage == MESA_SHADER_TESS_EVAL) {
3646 if (sel->nir->info.tess.point_mode)
3647 sel->rast_prim = MESA_PRIM_POINTS;
3648 else if (sel->nir->info.tess._primitive_mode == TESS_PRIMITIVE_ISOLINES)
3649 sel->rast_prim = MESA_PRIM_LINE_STRIP;
3650 else
3651 sel->rast_prim = MESA_PRIM_TRIANGLES;
3652 } else {
3653 sel->rast_prim = MESA_PRIM_TRIANGLES;
3654 }
3655 break;
3656 default:;
3657 }
3658
3659 bool ngg_culling_allowed =
3660 sscreen->info.gfx_level >= GFX10 &&
3661 sscreen->use_ngg_culling &&
3662 sel->info.writes_position &&
3663 !sel->info.writes_viewport_index && /* cull only against viewport 0 */
3664 !sel->nir->info.writes_memory &&
3665 /* NGG GS supports culling with streamout because it culls after streamout. */
3666 (sel->stage == MESA_SHADER_GEOMETRY || !sel->info.enabled_streamout_buffer_mask) &&
3667 (sel->stage != MESA_SHADER_GEOMETRY || sel->info.num_stream_output_components[0]) &&
3668 (sel->stage != MESA_SHADER_VERTEX ||
3669 (!sel->nir->info.vs.blit_sgprs_amd &&
3670 !sel->nir->info.vs.window_space_position));
3671
3672 sel->ngg_cull_vert_threshold = UINT_MAX; /* disabled (changed below) */
3673
3674 if (ngg_culling_allowed) {
3675 if (sel->stage == MESA_SHADER_VERTEX) {
3676 if (sscreen->debug_flags & DBG(ALWAYS_NGG_CULLING_ALL))
3677 sel->ngg_cull_vert_threshold = 0; /* always enabled */
3678 else
3679 sel->ngg_cull_vert_threshold = 128;
3680 } else if (sel->stage == MESA_SHADER_TESS_EVAL ||
3681 sel->stage == MESA_SHADER_GEOMETRY) {
3682 if (sel->rast_prim != MESA_PRIM_POINTS)
3683 sel->ngg_cull_vert_threshold = 0; /* always enabled */
3684 }
3685 }
3686
3687 (void)simple_mtx_init(&sel->mutex, mtx_plain);
3688
3689 si_schedule_initial_compile(sctx, sel->stage, &sel->ready, &sel->compiler_ctx_state,
3690 sel, si_init_shader_selector_async);
3691 return sel;
3692 }
3693
si_create_shader(struct pipe_context * ctx,const struct pipe_shader_state * state)3694 static void *si_create_shader(struct pipe_context *ctx, const struct pipe_shader_state *state)
3695 {
3696 struct si_context *sctx = (struct si_context *)ctx;
3697 struct si_screen *sscreen = (struct si_screen *)ctx->screen;
3698 bool cache_hit;
3699 struct si_shader_selector *sel = (struct si_shader_selector *)util_live_shader_cache_get(
3700 ctx, &sscreen->live_shader_cache, state, &cache_hit);
3701
3702 if (sel && cache_hit && sctx->debug.debug_message) {
3703 for (unsigned i = 0; i < 2; i++) {
3704 if (sel->main_shader_part[i])
3705 si_shader_dump_stats_for_shader_db(sscreen, sel->main_shader_part[i], &sctx->debug);
3706 if (sel->main_shader_part_ls[i])
3707 si_shader_dump_stats_for_shader_db(sscreen, sel->main_shader_part_ls[i], &sctx->debug);
3708 if (sel->main_shader_part_ngg[i])
3709 si_shader_dump_stats_for_shader_db(sscreen, sel->main_shader_part_ngg[i], &sctx->debug);
3710 if (sel->main_shader_part_ngg_es[i])
3711 si_shader_dump_stats_for_shader_db(sscreen, sel->main_shader_part_ngg_es[i], &sctx->debug);
3712 }
3713
3714 if (sel->main_shader_part_es)
3715 si_shader_dump_stats_for_shader_db(sscreen, sel->main_shader_part_es, &sctx->debug);
3716 }
3717 return sel;
3718 }
3719
si_update_streamout_state(struct si_context * sctx)3720 static void si_update_streamout_state(struct si_context *sctx)
3721 {
3722 struct si_shader_selector *shader_with_so = si_get_vs(sctx)->cso;
3723
3724 if (!shader_with_so)
3725 return;
3726
3727 sctx->streamout.enabled_stream_buffers_mask = shader_with_so->info.enabled_streamout_buffer_mask;
3728 sctx->streamout.stride_in_dw = shader_with_so->info.base.xfb_stride;
3729
3730 /* GDS must be allocated when any GDS instructions are used, otherwise it hangs. */
3731 if (sctx->gfx_level >= GFX11 && sctx->gfx_level < GFX12 &&
3732 shader_with_so->info.enabled_streamout_buffer_mask && !sctx->screen->gds_oa) {
3733 /* Gfx11 only uses GDS OA, not GDS memory. */
3734 simple_mtx_lock(&sctx->screen->gds_mutex);
3735 if (!sctx->screen->gds_oa) {
3736 sctx->screen->gds_oa = sctx->ws->buffer_create(sctx->ws, 1, 1, RADEON_DOMAIN_OA,
3737 RADEON_FLAG_DRIVER_INTERNAL);
3738 assert(sctx->screen->gds_oa);
3739 }
3740 simple_mtx_unlock(&sctx->screen->gds_mutex);
3741
3742 if (sctx->screen->gds_oa)
3743 sctx->ws->cs_add_buffer(&sctx->gfx_cs, sctx->screen->gds_oa, RADEON_USAGE_READWRITE,
3744 (enum radeon_bo_domain)0);
3745 }
3746 }
3747
si_update_clip_regs(struct si_context * sctx,struct si_shader_selector * old_hw_vs,struct si_shader * old_hw_vs_variant,struct si_shader_selector * next_hw_vs,struct si_shader * next_hw_vs_variant)3748 static void si_update_clip_regs(struct si_context *sctx, struct si_shader_selector *old_hw_vs,
3749 struct si_shader *old_hw_vs_variant,
3750 struct si_shader_selector *next_hw_vs,
3751 struct si_shader *next_hw_vs_variant)
3752 {
3753 if (next_hw_vs &&
3754 (!old_hw_vs ||
3755 (old_hw_vs->stage == MESA_SHADER_VERTEX && old_hw_vs->info.base.vs.window_space_position) !=
3756 (next_hw_vs->stage == MESA_SHADER_VERTEX && next_hw_vs->info.base.vs.window_space_position) ||
3757 old_hw_vs->info.clipdist_mask != next_hw_vs->info.clipdist_mask ||
3758 old_hw_vs->info.culldist_mask != next_hw_vs->info.culldist_mask || !old_hw_vs_variant ||
3759 !next_hw_vs_variant ||
3760 old_hw_vs_variant->pa_cl_vs_out_cntl != next_hw_vs_variant->pa_cl_vs_out_cntl))
3761 si_mark_atom_dirty(sctx, &sctx->atoms.s.clip_regs);
3762 }
3763
si_update_rasterized_prim(struct si_context * sctx)3764 static void si_update_rasterized_prim(struct si_context *sctx)
3765 {
3766 struct si_shader *hw_vs = si_get_vs(sctx)->current;
3767
3768 if (sctx->shader.gs.cso) {
3769 /* Only possibilities: POINTS, LINE_STRIP, TRIANGLES */
3770 si_set_rasterized_prim(sctx, sctx->shader.gs.cso->rast_prim, hw_vs, sctx->ngg);
3771 } else if (sctx->shader.tes.cso) {
3772 /* Only possibilities: POINTS, LINE_STRIP, TRIANGLES */
3773 si_set_rasterized_prim(sctx, sctx->shader.tes.cso->rast_prim, hw_vs, sctx->ngg);
3774 } else {
3775 /* The rasterized prim is determined by draw calls. */
3776 }
3777
3778 /* This must be done unconditionally because it also depends on si_shader fields. */
3779 si_update_ngg_sgpr_state_out_prim(sctx, hw_vs, sctx->ngg);
3780 }
3781
si_update_common_shader_state(struct si_context * sctx,struct si_shader_selector * sel,enum pipe_shader_type type)3782 static void si_update_common_shader_state(struct si_context *sctx, struct si_shader_selector *sel,
3783 enum pipe_shader_type type)
3784 {
3785 si_set_active_descriptors_for_shader(sctx, sel);
3786
3787 sctx->uses_bindless_samplers = si_shader_uses_bindless_samplers(sctx->shader.vs.cso) ||
3788 si_shader_uses_bindless_samplers(sctx->shader.gs.cso) ||
3789 si_shader_uses_bindless_samplers(sctx->shader.ps.cso) ||
3790 si_shader_uses_bindless_samplers(sctx->shader.tcs.cso) ||
3791 si_shader_uses_bindless_samplers(sctx->shader.tes.cso);
3792 sctx->uses_bindless_images = si_shader_uses_bindless_images(sctx->shader.vs.cso) ||
3793 si_shader_uses_bindless_images(sctx->shader.gs.cso) ||
3794 si_shader_uses_bindless_images(sctx->shader.ps.cso) ||
3795 si_shader_uses_bindless_images(sctx->shader.tcs.cso) ||
3796 si_shader_uses_bindless_images(sctx->shader.tes.cso);
3797
3798 if (type == PIPE_SHADER_VERTEX || type == PIPE_SHADER_TESS_EVAL || type == PIPE_SHADER_GEOMETRY)
3799 sctx->ngg_culling = 0; /* this will be enabled on the first draw if needed */
3800
3801 si_invalidate_inlinable_uniforms(sctx, type);
3802 sctx->do_update_shaders = true;
3803 }
3804
si_update_last_vgt_stage_state(struct si_context * sctx,struct si_shader_selector * old_hw_vs,struct si_shader * old_hw_vs_variant)3805 static void si_update_last_vgt_stage_state(struct si_context *sctx,
3806 /* hw_vs refers to the last VGT stage */
3807 struct si_shader_selector *old_hw_vs,
3808 struct si_shader *old_hw_vs_variant)
3809 {
3810 struct si_shader_ctx_state *hw_vs = si_get_vs(sctx);
3811
3812 si_update_vs_viewport_state(sctx);
3813 si_update_streamout_state(sctx);
3814 si_update_clip_regs(sctx, old_hw_vs, old_hw_vs_variant, hw_vs->cso, hw_vs->current);
3815 si_update_rasterized_prim(sctx);
3816
3817 /* Clear kill_pointsize because we only want it to be set in the last shader before PS. */
3818 sctx->shader.vs.key.ge.opt.kill_pointsize = 0;
3819 sctx->shader.tes.key.ge.opt.kill_pointsize = 0;
3820 sctx->shader.gs.key.ge.opt.kill_pointsize = 0;
3821 si_vs_ps_key_update_rast_prim_smooth_stipple(sctx);
3822 }
3823
si_bind_vs_shader(struct pipe_context * ctx,void * state)3824 static void si_bind_vs_shader(struct pipe_context *ctx, void *state)
3825 {
3826 struct si_context *sctx = (struct si_context *)ctx;
3827 struct si_shader_selector *old_hw_vs = si_get_vs(sctx)->cso;
3828 struct si_shader *old_hw_vs_variant = si_get_vs(sctx)->current;
3829 struct si_shader_selector *sel = (struct si_shader_selector*)state;
3830
3831 if (sctx->shader.vs.cso == sel)
3832 return;
3833
3834 sctx->shader.vs.cso = sel;
3835 sctx->shader.vs.current = (sel && sel->variants_count) ? sel->variants[0] : NULL;
3836 sctx->num_vs_blit_sgprs = sel ? sel->info.base.vs.blit_sgprs_amd : 0;
3837 sctx->vs_uses_draw_id = sel ? sel->info.uses_drawid : false;
3838
3839 if (si_update_ngg(sctx))
3840 si_shader_change_notify(sctx);
3841
3842 si_update_common_shader_state(sctx, sel, PIPE_SHADER_VERTEX);
3843 si_select_draw_vbo(sctx);
3844 si_update_last_vgt_stage_state(sctx, old_hw_vs, old_hw_vs_variant);
3845 si_vs_key_update_inputs(sctx);
3846
3847 if (sctx->screen->dpbb_allowed) {
3848 bool force_off = sel && sel->info.options & SI_PROFILE_VS_NO_BINNING;
3849
3850 if (force_off != sctx->dpbb_force_off_profile_vs) {
3851 sctx->dpbb_force_off_profile_vs = force_off;
3852 si_mark_atom_dirty(sctx, &sctx->atoms.s.dpbb_state);
3853 }
3854 }
3855 }
3856
si_update_tess_uses_prim_id(struct si_context * sctx)3857 static void si_update_tess_uses_prim_id(struct si_context *sctx)
3858 {
3859 sctx->ia_multi_vgt_param_key.u.tess_uses_prim_id =
3860 sctx->shader.tes.cso &&
3861 ((sctx->shader.tcs.cso && sctx->shader.tcs.cso->info.uses_primid) ||
3862 sctx->shader.tes.cso->info.uses_primid ||
3863 (sctx->shader.gs.cso && sctx->shader.gs.cso->info.uses_primid) ||
3864 (!sctx->shader.gs.cso && sctx->shader.ps.cso && sctx->shader.ps.cso->info.uses_primid));
3865 }
3866
si_update_ngg(struct si_context * sctx)3867 bool si_update_ngg(struct si_context *sctx)
3868 {
3869 if (!sctx->screen->use_ngg) {
3870 assert(!sctx->ngg);
3871 return false;
3872 }
3873
3874 bool new_ngg = true;
3875
3876 if (sctx->shader.gs.cso && sctx->shader.tes.cso && sctx->shader.gs.cso->tess_turns_off_ngg) {
3877 new_ngg = false;
3878 } else if (sctx->gfx_level < GFX11) {
3879 struct si_shader_selector *last = si_get_vs(sctx)->cso;
3880
3881 if ((last && last->info.enabled_streamout_buffer_mask) ||
3882 sctx->streamout.prims_gen_query_enabled)
3883 new_ngg = false;
3884 }
3885
3886 if (new_ngg != sctx->ngg) {
3887 /* Transitioning from NGG to legacy GS requires VGT_FLUSH on Navi10-14.
3888 * VGT_FLUSH is also emitted at the beginning of IBs when legacy GS ring
3889 * pointers are set.
3890 */
3891 if (sctx->screen->info.has_vgt_flush_ngg_legacy_bug && !new_ngg) {
3892 sctx->barrier_flags |= SI_BARRIER_EVENT_VGT_FLUSH;
3893 si_mark_atom_dirty(sctx, &sctx->atoms.s.barrier);
3894
3895 if (sctx->gfx_level == GFX10) {
3896 /* Workaround for https://gitlab.freedesktop.org/mesa/mesa/-/issues/2941 */
3897 si_flush_gfx_cs(sctx, RADEON_FLUSH_ASYNC_START_NEXT_GFX_IB_NOW, NULL);
3898 }
3899 }
3900
3901 sctx->ngg = new_ngg;
3902 si_select_draw_vbo(sctx);
3903 return true;
3904 }
3905 return false;
3906 }
3907
si_bind_gs_shader(struct pipe_context * ctx,void * state)3908 static void si_bind_gs_shader(struct pipe_context *ctx, void *state)
3909 {
3910 struct si_context *sctx = (struct si_context *)ctx;
3911 struct si_shader_selector *old_hw_vs = si_get_vs(sctx)->cso;
3912 struct si_shader *old_hw_vs_variant = si_get_vs(sctx)->current;
3913 struct si_shader_selector *sel = (struct si_shader_selector*)state;
3914 bool enable_changed = !!sctx->shader.gs.cso != !!sel;
3915 bool ngg_changed;
3916
3917 if (sctx->shader.gs.cso == sel)
3918 return;
3919
3920 sctx->shader.gs.cso = sel;
3921 sctx->shader.gs.current = (sel && sel->variants_count) ? sel->variants[0] : NULL;
3922 sctx->ia_multi_vgt_param_key.u.uses_gs = sel != NULL;
3923
3924 si_update_common_shader_state(sctx, sel, PIPE_SHADER_GEOMETRY);
3925 si_select_draw_vbo(sctx);
3926
3927 ngg_changed = si_update_ngg(sctx);
3928 if (ngg_changed || enable_changed)
3929 si_shader_change_notify(sctx);
3930 if (enable_changed) {
3931 if (sctx->ia_multi_vgt_param_key.u.uses_tess)
3932 si_update_tess_uses_prim_id(sctx);
3933 }
3934 si_update_last_vgt_stage_state(sctx, old_hw_vs, old_hw_vs_variant);
3935 }
3936
si_bind_tcs_shader(struct pipe_context * ctx,void * state)3937 static void si_bind_tcs_shader(struct pipe_context *ctx, void *state)
3938 {
3939 struct si_context *sctx = (struct si_context *)ctx;
3940 struct si_shader_selector *sel = (struct si_shader_selector*)state;
3941 bool enable_changed = !!sctx->shader.tcs.cso != !!sel;
3942
3943 /* Note it could happen that user shader sel is same as fixed function shader,
3944 * so we should update this field even sctx->shader.tcs.cso == sel.
3945 */
3946 sctx->is_user_tcs = !!sel;
3947
3948 if (sctx->shader.tcs.cso == sel)
3949 return;
3950
3951 sctx->shader.tcs.cso = sel;
3952 sctx->shader.tcs.current = (sel && sel->variants_count) ? sel->variants[0] : NULL;
3953 si_update_tess_uses_prim_id(sctx);
3954 si_update_tess_in_out_patch_vertices(sctx);
3955
3956 si_update_common_shader_state(sctx, sel, PIPE_SHADER_TESS_CTRL);
3957
3958 if (enable_changed)
3959 sctx->last_tcs = NULL; /* invalidate derived tess state */
3960 }
3961
si_bind_tes_shader(struct pipe_context * ctx,void * state)3962 static void si_bind_tes_shader(struct pipe_context *ctx, void *state)
3963 {
3964 struct si_context *sctx = (struct si_context *)ctx;
3965 struct si_shader_selector *old_hw_vs = si_get_vs(sctx)->cso;
3966 struct si_shader *old_hw_vs_variant = si_get_vs(sctx)->current;
3967 struct si_shader_selector *sel = (struct si_shader_selector*)state;
3968 bool enable_changed = !!sctx->shader.tes.cso != !!sel;
3969
3970 if (sctx->shader.tes.cso == sel)
3971 return;
3972
3973 sctx->shader.tes.cso = sel;
3974 sctx->shader.tes.current = (sel && sel->variants_count) ? sel->variants[0] : NULL;
3975 sctx->ia_multi_vgt_param_key.u.uses_tess = sel != NULL;
3976 si_update_tess_uses_prim_id(sctx);
3977
3978 sctx->shader.tcs.key.ge.opt.tes_prim_mode =
3979 sel ? sel->info.base.tess._primitive_mode : 0;
3980
3981 sctx->shader.tcs.key.ge.opt.tes_reads_tess_factors =
3982 sel ? sel->info.reads_tess_factors : 0;
3983
3984 if (sel) {
3985 sctx->tcs_offchip_layout &= 0x1fffffff;
3986 sctx->tcs_offchip_layout |=
3987 (sel->info.base.tess._primitive_mode << 29) |
3988 (sel->info.reads_tess_factors << 31);
3989
3990 si_mark_atom_dirty(sctx, &sctx->atoms.s.tess_io_layout);
3991 }
3992
3993 si_update_common_shader_state(sctx, sel, PIPE_SHADER_TESS_EVAL);
3994 si_select_draw_vbo(sctx);
3995
3996 bool ngg_changed = si_update_ngg(sctx);
3997 if (ngg_changed || enable_changed)
3998 si_shader_change_notify(sctx);
3999 if (enable_changed)
4000 sctx->last_tes_sh_base = -1; /* invalidate derived tess state */
4001 si_update_last_vgt_stage_state(sctx, old_hw_vs, old_hw_vs_variant);
4002 }
4003
si_update_vrs_flat_shading(struct si_context * sctx)4004 void si_update_vrs_flat_shading(struct si_context *sctx)
4005 {
4006 if (sctx->gfx_level >= GFX10_3 && sctx->shader.ps.cso) {
4007 struct si_state_rasterizer *rs = sctx->queued.named.rasterizer;
4008 struct si_shader_info *info = &sctx->shader.ps.cso->info;
4009 bool allow_flat_shading =
4010 info->allow_flat_shading && !sctx->framebuffer.disable_vrs_flat_shading &&
4011 !rs->line_smooth && !rs->poly_smooth && !rs->poly_stipple_enable &&
4012 !rs->point_smooth && (rs->flatshade || !info->uses_interp_color);
4013
4014 if (sctx->allow_flat_shading != allow_flat_shading) {
4015 sctx->allow_flat_shading = allow_flat_shading;
4016 si_mark_atom_dirty(sctx, &sctx->atoms.s.db_render_state);
4017 }
4018 }
4019 }
4020
si_bind_ps_shader(struct pipe_context * ctx,void * state)4021 static void si_bind_ps_shader(struct pipe_context *ctx, void *state)
4022 {
4023 struct si_context *sctx = (struct si_context *)ctx;
4024 struct si_shader_selector *old_sel = sctx->shader.ps.cso;
4025 struct si_shader_selector *sel = (struct si_shader_selector*)state;
4026
4027 /* skip if supplied shader is one already in use */
4028 if (old_sel == sel)
4029 return;
4030
4031 sctx->shader.ps.cso = sel;
4032 sctx->shader.ps.current = (sel && sel->variants_count) ? sel->variants[0] : NULL;
4033
4034 si_update_common_shader_state(sctx, sel, PIPE_SHADER_FRAGMENT);
4035 if (sel) {
4036 if (sctx->ia_multi_vgt_param_key.u.uses_tess)
4037 si_update_tess_uses_prim_id(sctx);
4038
4039 if (!old_sel || old_sel->info.colors_written != sel->info.colors_written)
4040 si_mark_atom_dirty(sctx, &sctx->atoms.s.cb_render_state);
4041
4042 if (sctx->screen->info.has_out_of_order_rast &&
4043 (!old_sel || old_sel->info.base.writes_memory != sel->info.base.writes_memory ||
4044 old_sel->info.base.fs.early_fragment_tests !=
4045 sel->info.base.fs.early_fragment_tests))
4046 si_mark_atom_dirty(sctx, &sctx->atoms.s.msaa_config);
4047 }
4048 si_update_ps_colorbuf0_slot(sctx);
4049
4050 si_ps_key_update_framebuffer(sctx);
4051 si_ps_key_update_framebuffer_blend_dsa_rasterizer(sctx);
4052 si_ps_key_update_rasterizer(sctx);
4053 si_ps_key_update_dsa(sctx);
4054 si_ps_key_update_sample_shading(sctx);
4055 si_ps_key_update_framebuffer_rasterizer_sample_shading(sctx);
4056 si_update_ps_inputs_read_or_disabled(sctx);
4057 si_update_vrs_flat_shading(sctx);
4058
4059 if (sctx->screen->dpbb_allowed) {
4060 bool force_off = sel && sel->info.options & SI_PROFILE_GFX9_GFX10_PS_NO_BINNING &&
4061 (sctx->gfx_level >= GFX9 && sctx->gfx_level <= GFX10_3);
4062
4063 if (force_off != sctx->dpbb_force_off_profile_ps) {
4064 sctx->dpbb_force_off_profile_ps = force_off;
4065 si_mark_atom_dirty(sctx, &sctx->atoms.s.dpbb_state);
4066 }
4067 }
4068 }
4069
si_delete_shader(struct si_context * sctx,struct si_shader * shader)4070 static void si_delete_shader(struct si_context *sctx, struct si_shader *shader)
4071 {
4072 if (shader->is_optimized) {
4073 util_queue_drop_job(&sctx->screen->shader_compiler_queue_opt_variants, &shader->ready);
4074 }
4075
4076 util_queue_fence_destroy(&shader->ready);
4077
4078 /* If destroyed shaders were not unbound, the next compiled
4079 * shader variant could get the same pointer address and so
4080 * binding it to the same shader stage would be considered
4081 * a no-op, causing random behavior.
4082 */
4083 int state_index = -1;
4084
4085 switch (shader->selector->stage) {
4086 case MESA_SHADER_VERTEX:
4087 if (shader->key.ge.as_ls) {
4088 if (sctx->gfx_level <= GFX8)
4089 state_index = SI_STATE_IDX(ls);
4090 } else if (shader->key.ge.as_es) {
4091 if (sctx->gfx_level <= GFX8)
4092 state_index = SI_STATE_IDX(es);
4093 } else if (shader->key.ge.as_ngg) {
4094 state_index = SI_STATE_IDX(gs);
4095 } else {
4096 state_index = SI_STATE_IDX(vs);
4097 }
4098 break;
4099 case MESA_SHADER_TESS_CTRL:
4100 state_index = SI_STATE_IDX(hs);
4101 break;
4102 case MESA_SHADER_TESS_EVAL:
4103 if (shader->key.ge.as_es) {
4104 if (sctx->gfx_level <= GFX8)
4105 state_index = SI_STATE_IDX(es);
4106 } else if (shader->key.ge.as_ngg) {
4107 state_index = SI_STATE_IDX(gs);
4108 } else {
4109 state_index = SI_STATE_IDX(vs);
4110 }
4111 break;
4112 case MESA_SHADER_GEOMETRY:
4113 if (shader->is_gs_copy_shader)
4114 state_index = SI_STATE_IDX(vs);
4115 else
4116 state_index = SI_STATE_IDX(gs);
4117 break;
4118 case MESA_SHADER_FRAGMENT:
4119 state_index = SI_STATE_IDX(ps);
4120 break;
4121 default:;
4122 }
4123
4124 if (shader->gs_copy_shader)
4125 si_delete_shader(sctx, shader->gs_copy_shader);
4126
4127 si_shader_selector_reference(sctx, &shader->previous_stage_sel, NULL);
4128 si_shader_destroy(shader);
4129 si_pm4_free_state(sctx, &shader->pm4, state_index);
4130 }
4131
si_destroy_shader_selector(struct pipe_context * ctx,void * cso)4132 static void si_destroy_shader_selector(struct pipe_context *ctx, void *cso)
4133 {
4134 struct si_context *sctx = (struct si_context *)ctx;
4135 struct si_shader_selector *sel = (struct si_shader_selector *)cso;
4136
4137 util_queue_drop_job(&sctx->screen->shader_compiler_queue, &sel->ready);
4138
4139 if (sctx->shaders[sel->stage].cso == sel) {
4140 sctx->shaders[sel->stage].cso = NULL;
4141 sctx->shaders[sel->stage].current = NULL;
4142 }
4143
4144 for (unsigned i = 0; i < sel->variants_count; i++) {
4145 si_delete_shader(sctx, sel->variants[i]);
4146 }
4147
4148 for (unsigned i = 0; i < 2; i++) {
4149 if (sel->main_shader_part[i])
4150 si_delete_shader(sctx, sel->main_shader_part[i]);
4151 if (sel->main_shader_part_ls[i])
4152 si_delete_shader(sctx, sel->main_shader_part_ls[i]);
4153 if (sel->main_shader_part_ngg[i])
4154 si_delete_shader(sctx, sel->main_shader_part_ngg[i]);
4155 if (sel->main_shader_part_ngg_es[i])
4156 si_delete_shader(sctx, sel->main_shader_part_ngg_es[i]);
4157 }
4158
4159 if (sel->main_shader_part_es)
4160 si_delete_shader(sctx, sel->main_shader_part_es);
4161
4162 free(sel->keys);
4163 free(sel->variants);
4164
4165 util_queue_fence_destroy(&sel->ready);
4166 simple_mtx_destroy(&sel->mutex);
4167 ralloc_free(sel->nir);
4168 free(sel->nir_binary);
4169 free(sel);
4170 }
4171
si_delete_shader_selector(struct pipe_context * ctx,void * state)4172 static void si_delete_shader_selector(struct pipe_context *ctx, void *state)
4173 {
4174 struct si_context *sctx = (struct si_context *)ctx;
4175 struct si_shader_selector *sel = (struct si_shader_selector *)state;
4176
4177 si_shader_selector_reference(sctx, &sel, NULL);
4178 }
4179
4180 /**
4181 * Writing CONFIG or UCONFIG VGT registers requires VGT_FLUSH before that.
4182 */
si_cs_preamble_add_vgt_flush(struct si_context * sctx,bool tmz)4183 static void si_cs_preamble_add_vgt_flush(struct si_context *sctx, bool tmz)
4184 {
4185 struct si_pm4_state *pm4 = tmz ? sctx->cs_preamble_state_tmz : sctx->cs_preamble_state;
4186 bool *has_vgt_flush = tmz ? &sctx->cs_preamble_has_vgt_flush_tmz :
4187 &sctx->cs_preamble_has_vgt_flush;
4188
4189 /* We shouldn't get here if registers are shadowed. */
4190 assert(!sctx->shadowing.registers);
4191
4192 if (*has_vgt_flush)
4193 return;
4194
4195 /* Done by Vulkan before VGT_FLUSH. */
4196 ac_pm4_cmd_add(&pm4->base, PKT3(PKT3_EVENT_WRITE, 0, 0));
4197 ac_pm4_cmd_add(&pm4->base, EVENT_TYPE(V_028A90_VS_PARTIAL_FLUSH) | EVENT_INDEX(4));
4198
4199 /* VGT_FLUSH is required even if VGT is idle. It resets VGT pointers. */
4200 ac_pm4_cmd_add(&pm4->base, PKT3(PKT3_EVENT_WRITE, 0, 0));
4201 ac_pm4_cmd_add(&pm4->base, EVENT_TYPE(V_028A90_VGT_FLUSH) | EVENT_INDEX(0));
4202 ac_pm4_finalize(&pm4->base);
4203
4204 *has_vgt_flush = true;
4205 }
4206
4207 /**
4208 * Writing CONFIG or UCONFIG VGT registers requires VGT_FLUSH before that.
4209 */
si_emit_vgt_flush(struct radeon_cmdbuf * cs)4210 static void si_emit_vgt_flush(struct radeon_cmdbuf *cs)
4211 {
4212 radeon_begin(cs);
4213
4214 /* This is required before VGT_FLUSH. */
4215 radeon_event_write(V_028A90_VS_PARTIAL_FLUSH);
4216
4217 /* VGT_FLUSH is required even if VGT is idle. It resets VGT pointers. */
4218 radeon_event_write(V_028A90_VGT_FLUSH);
4219 radeon_end();
4220 }
4221
4222 /* Initialize state related to ESGS / GSVS ring buffers */
si_update_gs_ring_buffers(struct si_context * sctx)4223 bool si_update_gs_ring_buffers(struct si_context *sctx)
4224 {
4225 assert(sctx->gfx_level < GFX11);
4226
4227 struct si_shader_selector *es =
4228 sctx->shader.tes.cso ? sctx->shader.tes.cso : sctx->shader.vs.cso;
4229 struct si_shader_selector *gs = sctx->shader.gs.cso;
4230
4231 /* Chip constants. */
4232 unsigned num_se = sctx->screen->info.max_se;
4233 unsigned wave_size = 64;
4234 unsigned max_gs_waves = 32 * num_se; /* max 32 per SE on GCN */
4235 /* On GFX6-GFX7, the value comes from VGT_GS_VERTEX_REUSE = 16.
4236 * On GFX8+, the value comes from VGT_VERTEX_REUSE_BLOCK_CNTL = 30 (+2).
4237 */
4238 unsigned gs_vertex_reuse = (sctx->gfx_level >= GFX8 ? 32 : 16) * num_se;
4239 unsigned alignment = 256 * num_se;
4240 /* The maximum size is 63.999 MB per SE. */
4241 unsigned max_size = ((unsigned)(63.999 * 1024 * 1024) & ~255) * num_se;
4242
4243 /* Calculate the minimum size. */
4244 unsigned min_esgs_ring_size = align(es->info.esgs_vertex_stride * gs_vertex_reuse * wave_size, alignment);
4245
4246 /* These are recommended sizes, not minimum sizes. */
4247 unsigned esgs_ring_size =
4248 max_gs_waves * 2 * wave_size * es->info.esgs_vertex_stride * gs->info.gs_input_verts_per_prim;
4249 unsigned gsvs_ring_size = max_gs_waves * 2 * wave_size * gs->info.max_gsvs_emit_size;
4250
4251 min_esgs_ring_size = align(min_esgs_ring_size, alignment);
4252 esgs_ring_size = align(esgs_ring_size, alignment);
4253 gsvs_ring_size = align(gsvs_ring_size, alignment);
4254
4255 esgs_ring_size = CLAMP(esgs_ring_size, min_esgs_ring_size, max_size);
4256 gsvs_ring_size = MIN2(gsvs_ring_size, max_size);
4257
4258 /* Some rings don't have to be allocated if shaders don't use them.
4259 * (e.g. no varyings between ES and GS or GS and VS)
4260 *
4261 * GFX9 doesn't have the ESGS ring.
4262 */
4263 bool update_esgs = sctx->gfx_level <= GFX8 && esgs_ring_size &&
4264 (!sctx->esgs_ring || sctx->esgs_ring->width0 < esgs_ring_size);
4265 bool update_gsvs =
4266 gsvs_ring_size && (!sctx->gsvs_ring || sctx->gsvs_ring->width0 < gsvs_ring_size);
4267
4268 if (!update_esgs && !update_gsvs)
4269 return true;
4270
4271 if (update_esgs) {
4272 pipe_resource_reference(&sctx->esgs_ring, NULL);
4273 sctx->esgs_ring =
4274 pipe_aligned_buffer_create(sctx->b.screen,
4275 PIPE_RESOURCE_FLAG_UNMAPPABLE | SI_RESOURCE_FLAG_DRIVER_INTERNAL |
4276 SI_RESOURCE_FLAG_DISCARDABLE,
4277 PIPE_USAGE_DEFAULT,
4278 esgs_ring_size, sctx->screen->info.pte_fragment_size);
4279 if (!sctx->esgs_ring)
4280 return false;
4281 }
4282
4283 if (update_gsvs) {
4284 pipe_resource_reference(&sctx->gsvs_ring, NULL);
4285 sctx->gsvs_ring =
4286 pipe_aligned_buffer_create(sctx->b.screen,
4287 PIPE_RESOURCE_FLAG_UNMAPPABLE | SI_RESOURCE_FLAG_DRIVER_INTERNAL |
4288 SI_RESOURCE_FLAG_DISCARDABLE,
4289 PIPE_USAGE_DEFAULT,
4290 gsvs_ring_size, sctx->screen->info.pte_fragment_size);
4291 if (!sctx->gsvs_ring)
4292 return false;
4293 }
4294
4295 /* Set ring bindings. */
4296 if (sctx->esgs_ring) {
4297 assert(sctx->gfx_level <= GFX8);
4298 si_set_ring_buffer(sctx, SI_RING_ESGS, sctx->esgs_ring, 0, sctx->esgs_ring->width0, false,
4299 false, 0, 0, 0);
4300 }
4301 if (sctx->gsvs_ring) {
4302 si_set_ring_buffer(sctx, SI_RING_GSVS, sctx->gsvs_ring, 0, sctx->gsvs_ring->width0, false,
4303 false, 0, 0, 0);
4304 }
4305
4306 if (sctx->shadowing.registers) {
4307 /* These registers will be shadowed, so set them only once. */
4308 struct radeon_cmdbuf *cs = &sctx->gfx_cs;
4309
4310 assert(sctx->gfx_level >= GFX7);
4311
4312 si_emit_vgt_flush(cs);
4313
4314 radeon_begin(cs);
4315
4316 /* Set the GS registers. */
4317 if (sctx->esgs_ring) {
4318 assert(sctx->gfx_level <= GFX8);
4319 radeon_set_uconfig_reg(R_030900_VGT_ESGS_RING_SIZE,
4320 sctx->esgs_ring->width0 / 256);
4321 }
4322 if (sctx->gsvs_ring) {
4323 radeon_set_uconfig_reg(R_030904_VGT_GSVS_RING_SIZE,
4324 sctx->gsvs_ring->width0 / 256);
4325 }
4326 radeon_end();
4327 return true;
4328 }
4329
4330 /* The codepath without register shadowing. */
4331 for (unsigned tmz = 0; tmz <= 1; tmz++) {
4332 struct si_pm4_state *pm4 = tmz ? sctx->cs_preamble_state_tmz : sctx->cs_preamble_state;
4333 uint16_t *gs_ring_state_dw_offset = tmz ? &sctx->gs_ring_state_dw_offset_tmz :
4334 &sctx->gs_ring_state_dw_offset;
4335 unsigned old_ndw = 0;
4336
4337 si_cs_preamble_add_vgt_flush(sctx, tmz);
4338
4339 if (!*gs_ring_state_dw_offset) {
4340 /* We are here for the first time. The packets will be added. */
4341 *gs_ring_state_dw_offset = pm4->base.ndw;
4342 } else {
4343 /* We have been here before. Overwrite the previous packets. */
4344 old_ndw = pm4->base.ndw;
4345 pm4->base.ndw = *gs_ring_state_dw_offset;
4346 }
4347
4348 /* Unallocated rings are written to reserve the space in the pm4
4349 * (to be able to overwrite them later). */
4350 if (sctx->gfx_level >= GFX7) {
4351 if (sctx->gfx_level <= GFX8)
4352 ac_pm4_set_reg(&pm4->base, R_030900_VGT_ESGS_RING_SIZE,
4353 sctx->esgs_ring ? sctx->esgs_ring->width0 / 256 : 0);
4354 ac_pm4_set_reg(&pm4->base, R_030904_VGT_GSVS_RING_SIZE,
4355 sctx->gsvs_ring ? sctx->gsvs_ring->width0 / 256 : 0);
4356 } else {
4357 ac_pm4_set_reg(&pm4->base, R_0088C8_VGT_ESGS_RING_SIZE,
4358 sctx->esgs_ring ? sctx->esgs_ring->width0 / 256 : 0);
4359 ac_pm4_set_reg(&pm4->base, R_0088CC_VGT_GSVS_RING_SIZE,
4360 sctx->gsvs_ring ? sctx->gsvs_ring->width0 / 256 : 0);
4361 }
4362 ac_pm4_finalize(&pm4->base);
4363
4364 if (old_ndw) {
4365 pm4->base.ndw = old_ndw;
4366 pm4->base.last_opcode = 255; /* invalid opcode (we don't save the last opcode) */
4367 }
4368 }
4369
4370 /* Flush the context to re-emit both cs_preamble states. */
4371 sctx->initial_gfx_cs_size = 0; /* force flush */
4372 si_flush_gfx_cs(sctx, RADEON_FLUSH_ASYNC_START_NEXT_GFX_IB_NOW, NULL);
4373
4374 return true;
4375 }
4376
si_shader_lock(struct si_shader * shader)4377 static void si_shader_lock(struct si_shader *shader)
4378 {
4379 simple_mtx_lock(&shader->selector->mutex);
4380 if (shader->previous_stage_sel) {
4381 assert(shader->previous_stage_sel != shader->selector);
4382 simple_mtx_lock(&shader->previous_stage_sel->mutex);
4383 }
4384 }
4385
si_shader_unlock(struct si_shader * shader)4386 static void si_shader_unlock(struct si_shader *shader)
4387 {
4388 if (shader->previous_stage_sel)
4389 simple_mtx_unlock(&shader->previous_stage_sel->mutex);
4390 simple_mtx_unlock(&shader->selector->mutex);
4391 }
4392
4393 /**
4394 * @returns 1 if \p sel has been updated to use a new scratch buffer
4395 * 0 if not
4396 * < 0 if there was a failure
4397 */
si_update_scratch_buffer(struct si_context * sctx,struct si_shader * shader)4398 static int si_update_scratch_buffer(struct si_context *sctx, struct si_shader *shader)
4399 {
4400 uint64_t scratch_va = sctx->scratch_buffer->gpu_address;
4401
4402 if (!shader)
4403 return 0;
4404
4405 /* This shader doesn't need a scratch buffer */
4406 if (shader->config.scratch_bytes_per_wave == 0)
4407 return 0;
4408
4409 /* Prevent race conditions when updating:
4410 * - si_shader::scratch_va
4411 * - si_shader::binary::code
4412 * - si_shader::previous_stage::binary::code.
4413 */
4414 si_shader_lock(shader);
4415
4416 /* This shader is already configured to use the current
4417 * scratch buffer. */
4418 if (shader->scratch_va == scratch_va) {
4419 si_shader_unlock(shader);
4420 return 0;
4421 }
4422
4423 assert(sctx->scratch_buffer);
4424
4425 /* Replace the shader bo with a new bo that has the relocs applied. */
4426 if (!si_shader_binary_upload(sctx->screen, shader, scratch_va)) {
4427 si_shader_unlock(shader);
4428 return -1;
4429 }
4430
4431 /* Update the shader state to use the new shader bo. */
4432 si_shader_init_pm4_state(sctx->screen, shader);
4433 shader->scratch_va = scratch_va;
4434
4435 si_shader_unlock(shader);
4436 return 1;
4437 }
4438
si_update_scratch_relocs(struct si_context * sctx)4439 static bool si_update_scratch_relocs(struct si_context *sctx)
4440 {
4441 int r;
4442
4443 /* Update the shaders, so that they are using the latest scratch.
4444 * The scratch buffer may have been changed since these shaders were
4445 * last used, so we still need to try to update them, even if they
4446 * require scratch buffers smaller than the current size.
4447 */
4448 r = si_update_scratch_buffer(sctx, sctx->shader.ps.current);
4449 if (r < 0)
4450 return false;
4451 if (r == 1)
4452 si_pm4_bind_state(sctx, ps, sctx->shader.ps.current);
4453
4454 r = si_update_scratch_buffer(sctx, sctx->shader.gs.current);
4455 if (r < 0)
4456 return false;
4457 if (r == 1)
4458 si_pm4_bind_state(sctx, gs, sctx->shader.gs.current);
4459
4460 r = si_update_scratch_buffer(sctx, sctx->shader.tcs.current);
4461 if (r < 0)
4462 return false;
4463 if (r == 1)
4464 si_pm4_bind_state(sctx, hs, sctx->shader.tcs.current);
4465
4466 /* VS can be bound as LS, ES, or VS. */
4467 r = si_update_scratch_buffer(sctx, sctx->shader.vs.current);
4468 if (r < 0)
4469 return false;
4470 if (r == 1) {
4471 if (sctx->shader.vs.current->key.ge.as_ls)
4472 si_pm4_bind_state(sctx, ls, sctx->shader.vs.current);
4473 else if (sctx->shader.vs.current->key.ge.as_es)
4474 si_pm4_bind_state(sctx, es, sctx->shader.vs.current);
4475 else if (sctx->shader.vs.current->key.ge.as_ngg)
4476 si_pm4_bind_state(sctx, gs, sctx->shader.vs.current);
4477 else
4478 si_pm4_bind_state(sctx, vs, sctx->shader.vs.current);
4479 }
4480
4481 /* TES can be bound as ES or VS. */
4482 r = si_update_scratch_buffer(sctx, sctx->shader.tes.current);
4483 if (r < 0)
4484 return false;
4485 if (r == 1) {
4486 if (sctx->shader.tes.current->key.ge.as_es)
4487 si_pm4_bind_state(sctx, es, sctx->shader.tes.current);
4488 else if (sctx->shader.tes.current->key.ge.as_ngg)
4489 si_pm4_bind_state(sctx, gs, sctx->shader.tes.current);
4490 else
4491 si_pm4_bind_state(sctx, vs, sctx->shader.tes.current);
4492 }
4493
4494 return true;
4495 }
4496
si_update_spi_tmpring_size(struct si_context * sctx,unsigned bytes)4497 bool si_update_spi_tmpring_size(struct si_context *sctx, unsigned bytes)
4498 {
4499 unsigned spi_tmpring_size;
4500 ac_get_scratch_tmpring_size(&sctx->screen->info, bytes,
4501 &sctx->max_seen_scratch_bytes_per_wave, &spi_tmpring_size);
4502
4503 unsigned scratch_needed_size = sctx->max_seen_scratch_bytes_per_wave *
4504 sctx->screen->info.max_scratch_waves;
4505
4506 if (scratch_needed_size > 0) {
4507 if (!sctx->scratch_buffer || scratch_needed_size > sctx->scratch_buffer->b.b.width0) {
4508 /* Create a bigger scratch buffer */
4509 si_resource_reference(&sctx->scratch_buffer, NULL);
4510
4511 sctx->scratch_buffer = si_aligned_buffer_create(
4512 &sctx->screen->b,
4513 PIPE_RESOURCE_FLAG_UNMAPPABLE | SI_RESOURCE_FLAG_DRIVER_INTERNAL |
4514 SI_RESOURCE_FLAG_DISCARDABLE,
4515 PIPE_USAGE_DEFAULT, scratch_needed_size,
4516 sctx->screen->info.pte_fragment_size);
4517 if (!sctx->scratch_buffer)
4518 return false;
4519 }
4520
4521 if (!sctx->screen->info.has_scratch_base_registers && !si_update_scratch_relocs(sctx))
4522 return false;
4523 }
4524
4525 if (spi_tmpring_size != sctx->spi_tmpring_size) {
4526 sctx->spi_tmpring_size = spi_tmpring_size;
4527 si_mark_atom_dirty(sctx, &sctx->atoms.s.scratch_state);
4528 }
4529 return true;
4530 }
4531
si_init_tess_factor_ring(struct si_context * sctx)4532 void si_init_tess_factor_ring(struct si_context *sctx)
4533 {
4534 struct si_screen *sscreen = sctx->screen;
4535 assert(!sctx->has_tessellation);
4536
4537 if (sctx->has_tessellation)
4538 return;
4539
4540 simple_mtx_lock(&sscreen->tess_ring_lock);
4541
4542 if (!sscreen->tess_rings) {
4543 /* The address must be aligned to 2^19, because the shader only
4544 * receives the high 13 bits. Align it to 2MB to match the GPU page size.
4545 */
4546 sscreen->tess_rings = pipe_aligned_buffer_create(sctx->b.screen,
4547 PIPE_RESOURCE_FLAG_UNMAPPABLE |
4548 SI_RESOURCE_FLAG_32BIT |
4549 SI_RESOURCE_FLAG_DRIVER_INTERNAL |
4550 SI_RESOURCE_FLAG_DISCARDABLE,
4551 PIPE_USAGE_DEFAULT,
4552 sscreen->hs.tess_offchip_ring_size +
4553 sscreen->hs.tess_factor_ring_size,
4554 2 * 1024 * 1024);
4555 if (!sscreen->tess_rings) {
4556 simple_mtx_unlock(&sscreen->tess_ring_lock);
4557 return;
4558 }
4559
4560 if (sscreen->info.has_tmz_support) {
4561 sscreen->tess_rings_tmz = pipe_aligned_buffer_create(sctx->b.screen,
4562 PIPE_RESOURCE_FLAG_UNMAPPABLE |
4563 PIPE_RESOURCE_FLAG_ENCRYPTED |
4564 SI_RESOURCE_FLAG_32BIT |
4565 SI_RESOURCE_FLAG_DRIVER_INTERNAL |
4566 SI_RESOURCE_FLAG_DISCARDABLE,
4567 PIPE_USAGE_DEFAULT,
4568 sscreen->hs.tess_offchip_ring_size +
4569 sscreen->hs.tess_factor_ring_size,
4570 2 * 1024 * 1024);
4571 }
4572 }
4573
4574 simple_mtx_unlock(&sscreen->tess_ring_lock);
4575 sctx->has_tessellation = true;
4576
4577 si_mark_atom_dirty(sctx, &sctx->atoms.s.spi_ge_ring_state);
4578 }
4579
si_emit_vgt_pipeline_state(struct si_context * sctx,unsigned index)4580 static void si_emit_vgt_pipeline_state(struct si_context *sctx, unsigned index)
4581 {
4582 struct radeon_cmdbuf *cs = &sctx->gfx_cs;
4583
4584 radeon_begin(cs);
4585 radeon_opt_set_context_reg(sctx->gfx_level >= GFX12 ?
4586 R_028A98_VGT_SHADER_STAGES_EN :
4587 R_028B54_VGT_SHADER_STAGES_EN,
4588 SI_TRACKED_VGT_SHADER_STAGES_EN, sctx->vgt_shader_stages_en);
4589 if (sctx->gfx_level == GFX10_3) {
4590 /* Legacy Tess+GS should disable reuse to prevent hangs on GFX10.3. */
4591 bool has_legacy_tess_gs = G_028B54_HS_EN(sctx->vgt_shader_stages_en) &&
4592 G_028B54_GS_EN(sctx->vgt_shader_stages_en) &&
4593 !G_028B54_PRIMGEN_EN(sctx->vgt_shader_stages_en); /* !NGG */
4594
4595 radeon_opt_set_context_reg(R_028AB4_VGT_REUSE_OFF, SI_TRACKED_VGT_REUSE_OFF,
4596 S_028AB4_REUSE_OFF(has_legacy_tess_gs));
4597 }
4598 radeon_end_update_context_roll();
4599
4600 if (sctx->gfx_level >= GFX10) {
4601 uint32_t ge_cntl = sctx->ge_cntl;
4602
4603 if (sctx->gfx_level < GFX11 && sctx->shader.tes.cso) {
4604 /* This must be a multiple of VGT_LS_HS_CONFIG.NUM_PATCHES. */
4605 ge_cntl |= S_03096C_PRIM_GRP_SIZE_GFX10(sctx->num_patches_per_workgroup);
4606 }
4607
4608 radeon_begin_again(cs);
4609 radeon_opt_set_uconfig_reg(R_03096C_GE_CNTL, SI_TRACKED_GE_CNTL, ge_cntl);
4610 radeon_end();
4611 }
4612 }
4613
si_emit_scratch_state(struct si_context * sctx,unsigned index)4614 static void si_emit_scratch_state(struct si_context *sctx, unsigned index)
4615 {
4616 struct radeon_cmdbuf *cs = &sctx->gfx_cs;
4617
4618 radeon_begin(cs);
4619 if (sctx->gfx_level >= GFX11) {
4620 radeon_set_context_reg_seq(R_0286E8_SPI_TMPRING_SIZE, 3);
4621 radeon_emit(sctx->spi_tmpring_size); /* SPI_TMPRING_SIZE */
4622 radeon_emit(sctx->scratch_buffer->gpu_address >> 8); /* SPI_GFX_SCRATCH_BASE_LO */
4623 radeon_emit(sctx->scratch_buffer->gpu_address >> 40); /* SPI_GFX_SCRATCH_BASE_HI */
4624 } else {
4625 radeon_set_context_reg(R_0286E8_SPI_TMPRING_SIZE, sctx->spi_tmpring_size);
4626 }
4627 radeon_end();
4628
4629 if (sctx->scratch_buffer) {
4630 radeon_add_to_buffer_list(sctx, &sctx->gfx_cs, sctx->scratch_buffer,
4631 RADEON_USAGE_READWRITE | RADEON_PRIO_SCRATCH_BUFFER);
4632 }
4633 }
4634
4635 struct si_fixed_func_tcs_shader_key {
4636 uint64_t outputs_written;
4637 uint8_t vertices_out;
4638 };
4639
4640 DERIVE_HASH_TABLE(si_fixed_func_tcs_shader_key);
4641
si_set_tcs_to_fixed_func_shader(struct si_context * sctx)4642 bool si_set_tcs_to_fixed_func_shader(struct si_context *sctx)
4643 {
4644 if (!sctx->fixed_func_tcs_shader_cache) {
4645 sctx->fixed_func_tcs_shader_cache = si_fixed_func_tcs_shader_key_table_create(NULL);
4646 }
4647
4648 struct si_fixed_func_tcs_shader_key key;
4649 key.outputs_written = sctx->shader.vs.cso->info.ls_es_outputs_written;
4650 key.vertices_out = sctx->patch_vertices;
4651
4652 struct hash_entry *entry = _mesa_hash_table_search(
4653 sctx->fixed_func_tcs_shader_cache, &key);
4654
4655 struct si_shader_selector *tcs;
4656 if (entry)
4657 tcs = (struct si_shader_selector *)entry->data;
4658 else {
4659 tcs = (struct si_shader_selector *)si_create_passthrough_tcs(sctx);
4660 if (!tcs)
4661 return false;
4662 _mesa_hash_table_insert(sctx->fixed_func_tcs_shader_cache, &key, (void *)tcs);
4663 }
4664
4665 sctx->shader.tcs.cso = tcs;
4666 return true;
4667 }
4668
si_update_tess_in_out_patch_vertices(struct si_context * sctx)4669 static void si_update_tess_in_out_patch_vertices(struct si_context *sctx)
4670 {
4671 if (sctx->is_user_tcs) {
4672 struct si_shader_selector *tcs = sctx->shader.tcs.cso;
4673
4674 bool same_patch_vertices =
4675 sctx->gfx_level >= GFX9 &&
4676 sctx->patch_vertices == tcs->info.base.tess.tcs_vertices_out;
4677
4678 if (sctx->shader.tcs.key.ge.opt.same_patch_vertices != same_patch_vertices) {
4679 sctx->shader.tcs.key.ge.opt.same_patch_vertices = same_patch_vertices;
4680 sctx->do_update_shaders = true;
4681 }
4682 } else {
4683 /* These fields are static for fixed function TCS. So no need to set
4684 * do_update_shaders between fixed-TCS draws. As fixed-TCS to user-TCS
4685 * or opposite, do_update_shaders should already be set by bind state.
4686 */
4687 sctx->shader.tcs.key.ge.opt.same_patch_vertices = sctx->gfx_level >= GFX9;
4688
4689 /* User may only change patch vertices, needs to update fixed func TCS. */
4690 if (sctx->shader.tcs.cso &&
4691 sctx->shader.tcs.cso->info.base.tess.tcs_vertices_out != sctx->patch_vertices)
4692 sctx->do_update_shaders = true;
4693 }
4694 }
4695
si_set_patch_vertices(struct pipe_context * ctx,uint8_t patch_vertices)4696 static void si_set_patch_vertices(struct pipe_context *ctx, uint8_t patch_vertices)
4697 {
4698 struct si_context *sctx = (struct si_context *)ctx;
4699
4700 if (sctx->patch_vertices != patch_vertices) {
4701 sctx->patch_vertices = patch_vertices;
4702 si_update_tess_in_out_patch_vertices(sctx);
4703 if (sctx->shader.tcs.current) {
4704 /* Update the io layout now if possible,
4705 * otherwise make sure it's done by si_update_shaders.
4706 */
4707 if (sctx->has_tessellation)
4708 si_update_tess_io_layout_state(sctx);
4709 else
4710 sctx->do_update_shaders = true;
4711 }
4712
4713 /* Gfx12 programs patch_vertices in VGT_PRIMITIVE_TYPE.NUM_INPUT_CP. Make sure
4714 * the register is updated.
4715 */
4716 if (sctx->gfx_level >= GFX12 && sctx->last_prim == MESA_PRIM_PATCHES)
4717 sctx->last_prim = -1;
4718 }
4719 }
4720
si_shader_lshs_vertex_stride(struct si_shader * ls)4721 unsigned si_shader_lshs_vertex_stride(struct si_shader *ls)
4722 {
4723 unsigned num_slots;
4724
4725 if (ls->selector->stage == MESA_SHADER_VERTEX && !ls->next_shader) {
4726 assert(ls->key.ge.as_ls);
4727 assert(ls->selector->screen->info.gfx_level <= GFX8 || !ls->is_monolithic);
4728 num_slots = util_last_bit64(ls->selector->info.ls_es_outputs_written);
4729 } else {
4730 struct si_shader *tcs = ls->next_shader ? ls->next_shader : ls;
4731
4732 assert(tcs->selector->stage == MESA_SHADER_TESS_CTRL);
4733 assert(tcs->selector->screen->info.gfx_level >= GFX9);
4734
4735 if (tcs->is_monolithic) {
4736 uint64_t lds_inputs_read = tcs->selector->info.tcs_inputs_via_lds;
4737
4738 /* If the TCS in/out number of vertices is different, all inputs are passed via LDS. */
4739 if (!tcs->key.ge.opt.same_patch_vertices)
4740 lds_inputs_read |= tcs->selector->info.tcs_inputs_via_temp;
4741
4742 /* NIR lowering passes pack LS outputs/HS inputs if the usage masks of both are known. */
4743 num_slots = util_bitcount64(lds_inputs_read);
4744 } else {
4745 num_slots = util_last_bit64(tcs->previous_stage_sel->info.ls_es_outputs_written);
4746 }
4747 }
4748
4749 /* Add 1 dword to reduce LDS bank conflicts, so that each vertex starts on a different LDS
4750 * bank.
4751 */
4752 return num_slots ? num_slots * 16 + 4 : 0;
4753 }
4754
4755 /**
4756 * This calculates the LDS size for tessellation shaders (VS, TCS, TES).
4757 * LS.LDS_SIZE is shared by all 3 shader stages.
4758 *
4759 * The information about LDS and other non-compile-time parameters is then
4760 * written to userdata SGPRs.
4761 *
4762 * This depends on:
4763 * - patch_vertices
4764 * - VS and the currently selected shader variant (called by si_update_shaders)
4765 * - TCS and the currently selected shader variant (called by si_update_shaders)
4766 * - tess_uses_prim_id (called by si_update_shaders)
4767 * - sh_base[TESS_EVAL] depending on GS on/off (called by si_update_shaders)
4768 */
si_update_tess_io_layout_state(struct si_context * sctx)4769 void si_update_tess_io_layout_state(struct si_context *sctx)
4770 {
4771 struct si_shader *ls_current;
4772 struct si_shader_selector *tcs = sctx->shader.tcs.cso;
4773 bool tess_uses_primid = sctx->ia_multi_vgt_param_key.u.tess_uses_prim_id;
4774 bool has_primid_instancing_bug = sctx->gfx_level == GFX6 && sctx->screen->info.max_se == 1;
4775 unsigned tes_sh_base = sctx->shader_pointers.sh_base[PIPE_SHADER_TESS_EVAL];
4776 uint8_t num_tcs_input_cp = sctx->patch_vertices;
4777
4778 assert(sctx->shader.tcs.current);
4779
4780 /* Since GFX9 has merged LS-HS in the TCS state, set LS = TCS. */
4781 if (sctx->gfx_level >= GFX9) {
4782 ls_current = sctx->shader.tcs.current;
4783 } else {
4784 ls_current = sctx->shader.vs.current;
4785
4786 if (!ls_current) {
4787 sctx->do_update_shaders = true;
4788 return;
4789 }
4790 }
4791
4792 if (sctx->last_ls == ls_current && sctx->last_tcs == tcs &&
4793 sctx->last_tes_sh_base == tes_sh_base && sctx->last_num_tcs_input_cp == num_tcs_input_cp &&
4794 (!has_primid_instancing_bug || (sctx->last_tess_uses_primid == tess_uses_primid)))
4795 return;
4796
4797 sctx->last_ls = ls_current;
4798 sctx->last_tcs = tcs;
4799 sctx->last_tes_sh_base = tes_sh_base;
4800 sctx->last_num_tcs_input_cp = num_tcs_input_cp;
4801 sctx->last_tess_uses_primid = tess_uses_primid;
4802
4803 /* This calculates how shader inputs and outputs among VS, TCS, and TES
4804 * are laid out in LDS and memory.
4805 */
4806 unsigned num_tcs_output_cp = tcs->info.base.tess.tcs_vertices_out;
4807 unsigned lds_input_vertex_size = si_shader_lshs_vertex_stride(ls_current);
4808 unsigned num_mem_tcs_outputs = util_last_bit64(tcs->info.tcs_outputs_written_for_tes);
4809 unsigned num_mem_tcs_patch_outputs =
4810 util_last_bit(tcs->info.patch_outputs_written_for_tes |
4811 (!ls_current->is_monolithic || ls_current->key.ge.opt.tes_reads_tess_factors ?
4812 tcs->info.tess_levels_written_for_tes : 0));
4813 unsigned num_patches, lds_size;
4814
4815 /* Compute NUM_PATCHES and LDS_SIZE. */
4816 ac_nir_compute_tess_wg_info(&sctx->screen->info, &tcs->info.base, ls_current->wave_size,
4817 tess_uses_primid, tcs->info.tessfactors_are_def_in_all_invocs,
4818 num_tcs_input_cp, lds_input_vertex_size,
4819 num_mem_tcs_outputs, num_mem_tcs_patch_outputs,
4820 &num_patches, &lds_size);
4821
4822 if (sctx->num_patches_per_workgroup != num_patches) {
4823 sctx->num_patches_per_workgroup = num_patches;
4824 si_mark_atom_dirty(sctx, &sctx->atoms.s.vgt_pipeline_state);
4825 }
4826
4827 /* Compute userdata SGPRs. */
4828 unsigned num_lds_vs_outputs = lds_input_vertex_size / 16;
4829 assert(ls_current->config.lds_size == 0);
4830 assert(num_tcs_input_cp <= 32);
4831 assert(num_tcs_output_cp <= 32);
4832 assert(num_patches <= 128);
4833 assert(num_lds_vs_outputs <= 63);
4834 assert(num_mem_tcs_outputs <= 63);
4835
4836 uint64_t ring_va =
4837 sctx->ws->cs_is_secure(&sctx->gfx_cs) ?
4838 si_resource(sctx->screen->tess_rings_tmz)->gpu_address :
4839 si_resource(sctx->screen->tess_rings)->gpu_address;
4840 assert((ring_va & u_bit_consecutive(0, 19)) == 0);
4841
4842 sctx->tes_offchip_ring_va_sgpr = ring_va;
4843 sctx->tcs_offchip_layout &= 0xe0000000;
4844 sctx->tcs_offchip_layout |=
4845 (num_patches - 1) | ((num_tcs_output_cp - 1) << 7) | ((num_tcs_input_cp - 1) << 12) |
4846 (num_lds_vs_outputs << 17) | (num_mem_tcs_outputs << 23);
4847
4848 unsigned ls_hs_rsrc2;
4849
4850 if (sctx->gfx_level >= GFX9) {
4851 ls_hs_rsrc2 = sctx->shader.tcs.current->config.rsrc2;
4852
4853 if (sctx->gfx_level >= GFX10)
4854 ls_hs_rsrc2 |= S_00B42C_LDS_SIZE_GFX10(lds_size);
4855 else
4856 ls_hs_rsrc2 |= S_00B42C_LDS_SIZE_GFX9(lds_size);
4857 } else {
4858 ls_hs_rsrc2 = sctx->shader.vs.current->config.rsrc2;
4859
4860 si_multiwave_lds_size_workaround(sctx->screen, &lds_size);
4861 ls_hs_rsrc2 |= S_00B52C_LDS_SIZE(lds_size);
4862 }
4863
4864 sctx->ls_hs_rsrc2 = ls_hs_rsrc2;
4865 sctx->ls_hs_config =
4866 S_028B58_NUM_PATCHES(sctx->num_patches_per_workgroup) |
4867 S_028B58_HS_NUM_OUTPUT_CP(num_tcs_output_cp);
4868
4869 if (sctx->gfx_level < GFX12)
4870 sctx->ls_hs_config |= S_028B58_HS_NUM_INPUT_CP(num_tcs_input_cp);
4871
4872 si_mark_atom_dirty(sctx, &sctx->atoms.s.tess_io_layout);
4873 }
4874
gfx6_emit_tess_io_layout_state(struct si_context * sctx,unsigned index)4875 static void gfx6_emit_tess_io_layout_state(struct si_context *sctx, unsigned index)
4876 {
4877 struct radeon_cmdbuf *cs = &sctx->gfx_cs;
4878
4879 if (!sctx->shader.tes.cso || !sctx->shader.tcs.current)
4880 return;
4881
4882 radeon_begin(cs);
4883 if (sctx->gfx_level >= GFX12) {
4884 gfx12_opt_push_gfx_sh_reg(R_00B42C_SPI_SHADER_PGM_RSRC2_HS,
4885 SI_TRACKED_SPI_SHADER_PGM_RSRC2_HS, sctx->ls_hs_rsrc2);
4886
4887 /* Set userdata SGPRs for merged LS-HS. */
4888 gfx12_opt_push_gfx_sh_reg(R_00B430_SPI_SHADER_USER_DATA_HS_0 +
4889 GFX9_SGPR_TCS_OFFCHIP_LAYOUT * 4,
4890 SI_TRACKED_SPI_SHADER_USER_DATA_HS__TCS_OFFCHIP_LAYOUT,
4891 sctx->tcs_offchip_layout);
4892 gfx12_opt_push_gfx_sh_reg(R_00B430_SPI_SHADER_USER_DATA_HS_0 +
4893 GFX9_SGPR_TCS_OFFCHIP_ADDR * 4,
4894 SI_TRACKED_SPI_SHADER_USER_DATA_HS__TCS_OFFCHIP_ADDR,
4895 sctx->tes_offchip_ring_va_sgpr);
4896 } else if (sctx->screen->info.has_set_sh_pairs_packed) {
4897 gfx11_opt_push_gfx_sh_reg(R_00B42C_SPI_SHADER_PGM_RSRC2_HS,
4898 SI_TRACKED_SPI_SHADER_PGM_RSRC2_HS, sctx->ls_hs_rsrc2);
4899
4900 /* Set userdata SGPRs for merged LS-HS. */
4901 gfx11_opt_push_gfx_sh_reg(R_00B430_SPI_SHADER_USER_DATA_HS_0 +
4902 GFX9_SGPR_TCS_OFFCHIP_LAYOUT * 4,
4903 SI_TRACKED_SPI_SHADER_USER_DATA_HS__TCS_OFFCHIP_LAYOUT,
4904 sctx->tcs_offchip_layout);
4905 gfx11_opt_push_gfx_sh_reg(R_00B430_SPI_SHADER_USER_DATA_HS_0 +
4906 GFX9_SGPR_TCS_OFFCHIP_ADDR * 4,
4907 SI_TRACKED_SPI_SHADER_USER_DATA_HS__TCS_OFFCHIP_ADDR,
4908 sctx->tes_offchip_ring_va_sgpr);
4909 } else if (sctx->gfx_level >= GFX9) {
4910 radeon_opt_set_sh_reg(R_00B42C_SPI_SHADER_PGM_RSRC2_HS,
4911 SI_TRACKED_SPI_SHADER_PGM_RSRC2_HS, sctx->ls_hs_rsrc2);
4912
4913 /* Set userdata SGPRs for merged LS-HS. */
4914 radeon_opt_set_sh_reg2(R_00B430_SPI_SHADER_USER_DATA_HS_0 +
4915 GFX9_SGPR_TCS_OFFCHIP_LAYOUT * 4,
4916 SI_TRACKED_SPI_SHADER_USER_DATA_HS__TCS_OFFCHIP_LAYOUT,
4917 sctx->tcs_offchip_layout, sctx->tes_offchip_ring_va_sgpr);
4918 } else {
4919 /* Due to a hw bug, RSRC2_LS must be written twice with another
4920 * LS register written in between. */
4921 if (sctx->gfx_level == GFX7 && sctx->family != CHIP_HAWAII)
4922 radeon_set_sh_reg(R_00B52C_SPI_SHADER_PGM_RSRC2_LS, sctx->ls_hs_rsrc2);
4923 radeon_set_sh_reg_seq(R_00B528_SPI_SHADER_PGM_RSRC1_LS, 2);
4924 radeon_emit(sctx->shader.vs.current->config.rsrc1);
4925 radeon_emit(sctx->ls_hs_rsrc2);
4926
4927 /* Set userdata SGPRs for TCS. */
4928 radeon_opt_set_sh_reg3(R_00B430_SPI_SHADER_USER_DATA_HS_0 +
4929 GFX6_SGPR_TCS_OFFCHIP_LAYOUT * 4,
4930 SI_TRACKED_SPI_SHADER_USER_DATA_HS__TCS_OFFCHIP_LAYOUT,
4931 sctx->tcs_offchip_layout, sctx->tes_offchip_ring_va_sgpr,
4932 sctx->current_vs_state);
4933 }
4934
4935 /* Set userdata SGPRs for TES. */
4936 unsigned tes_sh_base = sctx->shader_pointers.sh_base[PIPE_SHADER_TESS_EVAL];
4937 assert(tes_sh_base);
4938
4939 /* TES (as ES or VS) reuses the BaseVertex and DrawID user SGPRs that are used when
4940 * tessellation is disabled. We can do that because those user SGPRs are only set in LS
4941 * for tessellation and are unused in TES.
4942 */
4943 if (sctx->screen->info.has_set_sh_pairs_packed) {
4944 gfx11_opt_push_gfx_sh_reg(tes_sh_base + SI_SGPR_TES_OFFCHIP_LAYOUT * 4,
4945 SI_TRACKED_SPI_SHADER_USER_DATA_ES__BASE_VERTEX,
4946 sctx->tcs_offchip_layout);
4947 gfx11_opt_push_gfx_sh_reg(tes_sh_base + SI_SGPR_TES_OFFCHIP_ADDR * 4,
4948 SI_TRACKED_SPI_SHADER_USER_DATA_ES__DRAWID,
4949 sctx->tes_offchip_ring_va_sgpr);
4950 } else if (sctx->ngg || sctx->shader.gs.cso) {
4951 radeon_opt_set_sh_reg2(tes_sh_base + SI_SGPR_TES_OFFCHIP_LAYOUT * 4,
4952 SI_TRACKED_SPI_SHADER_USER_DATA_ES__BASE_VERTEX,
4953 sctx->tcs_offchip_layout, sctx->tes_offchip_ring_va_sgpr);
4954 } else {
4955 radeon_opt_set_sh_reg2(tes_sh_base + SI_SGPR_TES_OFFCHIP_LAYOUT * 4,
4956 SI_TRACKED_SPI_SHADER_USER_DATA_VS__BASE_VERTEX,
4957 sctx->tcs_offchip_layout, sctx->tes_offchip_ring_va_sgpr);
4958 }
4959 radeon_end();
4960
4961 radeon_begin_again(cs);
4962 if (sctx->gfx_level >= GFX7) {
4963 radeon_opt_set_context_reg_idx(R_028B58_VGT_LS_HS_CONFIG,
4964 SI_TRACKED_VGT_LS_HS_CONFIG, 2, sctx->ls_hs_config);
4965 } else {
4966 radeon_opt_set_context_reg(R_028B58_VGT_LS_HS_CONFIG,
4967 SI_TRACKED_VGT_LS_HS_CONFIG, sctx->ls_hs_config);
4968 }
4969 radeon_end_update_context_roll();
4970 }
4971
gfx12_emit_tess_io_layout_state(struct si_context * sctx,unsigned index)4972 static void gfx12_emit_tess_io_layout_state(struct si_context *sctx, unsigned index)
4973 {
4974 struct radeon_cmdbuf *cs = &sctx->gfx_cs;
4975
4976 if (!sctx->shader.tes.cso || !sctx->shader.tcs.current)
4977 return;
4978
4979 gfx12_opt_push_gfx_sh_reg(R_00B42C_SPI_SHADER_PGM_RSRC2_HS,
4980 SI_TRACKED_SPI_SHADER_PGM_RSRC2_HS, sctx->ls_hs_rsrc2);
4981 /* Set userdata SGPRs for merged LS-HS. */
4982 gfx12_opt_push_gfx_sh_reg(R_00B430_SPI_SHADER_USER_DATA_HS_0 +
4983 GFX9_SGPR_TCS_OFFCHIP_LAYOUT * 4,
4984 SI_TRACKED_SPI_SHADER_USER_DATA_HS__TCS_OFFCHIP_LAYOUT,
4985 sctx->tcs_offchip_layout);
4986 gfx12_opt_push_gfx_sh_reg(R_00B430_SPI_SHADER_USER_DATA_HS_0 +
4987 GFX9_SGPR_TCS_OFFCHIP_ADDR * 4,
4988 SI_TRACKED_SPI_SHADER_USER_DATA_HS__TCS_OFFCHIP_ADDR,
4989 sctx->tes_offchip_ring_va_sgpr);
4990
4991 /* Set userdata SGPRs for TES. */
4992 unsigned tes_sh_base = sctx->shader_pointers.sh_base[PIPE_SHADER_TESS_EVAL];
4993 assert(tes_sh_base);
4994
4995 /* TES (as ES or VS) reuses the BaseVertex and DrawID user SGPRs that are used when
4996 * tessellation is disabled. We can do that because those user SGPRs are only set in LS
4997 * for tessellation and are unused in TES.
4998 */
4999 gfx12_opt_push_gfx_sh_reg(tes_sh_base + SI_SGPR_TES_OFFCHIP_LAYOUT * 4,
5000 SI_TRACKED_SPI_SHADER_USER_DATA_ES__BASE_VERTEX,
5001 sctx->tcs_offchip_layout);
5002 gfx12_opt_push_gfx_sh_reg(tes_sh_base + SI_SGPR_TES_OFFCHIP_ADDR * 4,
5003 SI_TRACKED_SPI_SHADER_USER_DATA_ES__DRAWID,
5004 sctx->tes_offchip_ring_va_sgpr);
5005
5006 radeon_begin(cs);
5007 radeon_opt_set_context_reg_idx(R_028B58_VGT_LS_HS_CONFIG,
5008 SI_TRACKED_VGT_LS_HS_CONFIG, 2, sctx->ls_hs_config);
5009 radeon_end(); /* don't track context rolls on GFX12 */
5010 }
5011
si_init_screen_live_shader_cache(struct si_screen * sscreen)5012 void si_init_screen_live_shader_cache(struct si_screen *sscreen)
5013 {
5014 util_live_shader_cache_init(&sscreen->live_shader_cache, si_create_shader_selector,
5015 si_destroy_shader_selector);
5016 }
5017
5018 template<int NUM_INTERP>
si_emit_spi_map(struct si_context * sctx,unsigned index)5019 static void si_emit_spi_map(struct si_context *sctx, unsigned index)
5020 {
5021 struct si_shader *ps = sctx->shader.ps.current;
5022 struct si_shader *vs = si_get_vs(sctx)->current;
5023 unsigned spi_ps_input_cntl[NUM_INTERP];
5024
5025 STATIC_ASSERT(NUM_INTERP >= 0 && NUM_INTERP <= 32);
5026
5027 if (sctx->gfx_level >= GFX12) {
5028 gfx12_opt_push_gfx_sh_reg(R_00B0C4_SPI_SHADER_GS_OUT_CONFIG_PS,
5029 SI_TRACKED_SPI_SHADER_GS_OUT_CONFIG_PS,
5030 vs->ngg.spi_vs_out_config | ps->ps.spi_gs_out_config_ps);
5031 }
5032
5033 if (!NUM_INTERP)
5034 return;
5035
5036 struct si_state_rasterizer *rs = sctx->queued.named.rasterizer;
5037
5038 for (unsigned i = 0; i < NUM_INTERP; i++) {
5039 union si_ps_input_info input = ps->info.ps_inputs[i];
5040 unsigned ps_input_cntl = vs->info.vs_output_ps_input_cntl[input.semantic];
5041 bool non_default_val = G_028644_OFFSET(ps_input_cntl) != 0x20;
5042
5043 if (non_default_val) {
5044 if (input.interpolate == INTERP_MODE_FLAT ||
5045 (input.interpolate == INTERP_MODE_COLOR && rs->flatshade))
5046 ps_input_cntl |= S_028644_FLAT_SHADE(1);
5047
5048 if (input.fp16_lo_hi_valid) {
5049 ps_input_cntl |= S_028644_FP16_INTERP_MODE(1) |
5050 S_028644_ATTR0_VALID(1) | /* this must be set if FP16_INTERP_MODE is set */
5051 S_028644_ATTR1_VALID(!!(input.fp16_lo_hi_valid & 0x2));
5052 }
5053 }
5054
5055 if (input.semantic == VARYING_SLOT_PNTC ||
5056 (input.semantic >= VARYING_SLOT_TEX0 && input.semantic <= VARYING_SLOT_TEX7 &&
5057 rs->sprite_coord_enable & (1 << (input.semantic - VARYING_SLOT_TEX0)))) {
5058 /* Overwrite the whole value (except OFFSET) for sprite coordinates. */
5059 ps_input_cntl &= ~C_028644_OFFSET;
5060 ps_input_cntl |= S_028644_PT_SPRITE_TEX(1);
5061 if (input.fp16_lo_hi_valid & 0x1) {
5062 ps_input_cntl |= S_028644_FP16_INTERP_MODE(1) |
5063 S_028644_ATTR0_VALID(1);
5064 }
5065 }
5066
5067 spi_ps_input_cntl[i] = ps_input_cntl;
5068 }
5069
5070 /* Performance notes:
5071 * Dota 2: Only ~16% of SPI map updates set different values.
5072 * Talos: Only ~9% of SPI map updates set different values.
5073 */
5074 if (sctx->gfx_level >= GFX12) {
5075 radeon_begin(&sctx->gfx_cs);
5076 radeon_opt_set_context_regn(R_028664_SPI_PS_INPUT_CNTL_0, spi_ps_input_cntl,
5077 sctx->tracked_regs.spi_ps_input_cntl, NUM_INTERP);
5078 radeon_end(); /* don't track context rolls on GFX12 */
5079 } else {
5080 radeon_begin(&sctx->gfx_cs);
5081 radeon_opt_set_context_regn(R_028644_SPI_PS_INPUT_CNTL_0, spi_ps_input_cntl,
5082 sctx->tracked_regs.spi_ps_input_cntl, NUM_INTERP);
5083 radeon_end_update_context_roll();
5084 }
5085 }
5086
si_emit_spi_ge_ring_state(struct si_context * sctx,unsigned index)5087 static void si_emit_spi_ge_ring_state(struct si_context *sctx, unsigned index)
5088 {
5089 struct si_screen *sscreen = sctx->screen;
5090
5091 if (sctx->has_tessellation) {
5092 struct pipe_resource *tf_ring =
5093 sctx->ws->cs_is_secure(&sctx->gfx_cs) ? sscreen->tess_rings_tmz : sscreen->tess_rings;
5094 uint64_t factor_va = si_resource(tf_ring)->gpu_address +
5095 sscreen->hs.tess_offchip_ring_size;
5096
5097 unsigned tf_ring_size_field = sscreen->hs.tess_factor_ring_size / 4;
5098 if (sctx->gfx_level >= GFX11)
5099 tf_ring_size_field /= sscreen->info.max_se;
5100
5101 assert((tf_ring_size_field & C_030938_SIZE) == 0);
5102
5103 radeon_add_to_buffer_list(sctx, &sctx->gfx_cs, si_resource(tf_ring),
5104 RADEON_USAGE_READWRITE | RADEON_PRIO_SHADER_RINGS);
5105
5106 radeon_begin(&sctx->gfx_cs);
5107 /* Required before writing tessellation config registers. */
5108 radeon_event_write(V_028A90_VS_PARTIAL_FLUSH);
5109 radeon_event_write(V_028A90_VGT_FLUSH);
5110
5111 if (sctx->gfx_level >= GFX7) {
5112 radeon_set_uconfig_reg_seq(R_030938_VGT_TF_RING_SIZE, 3);
5113 radeon_emit(S_030938_SIZE(tf_ring_size_field)); /* R_030938_VGT_TF_RING_SIZE */
5114 radeon_emit(sscreen->hs.hs_offchip_param); /* R_03093C_VGT_HS_OFFCHIP_PARAM */
5115 radeon_emit(factor_va >> 8); /* R_030940_VGT_TF_MEMORY_BASE */
5116
5117 if (sctx->gfx_level >= GFX12)
5118 radeon_set_uconfig_reg(R_03099C_VGT_TF_MEMORY_BASE_HI, S_03099C_BASE_HI(factor_va >> 40));
5119 else if (sctx->gfx_level >= GFX10)
5120 radeon_set_uconfig_reg(R_030984_VGT_TF_MEMORY_BASE_HI, S_030984_BASE_HI(factor_va >> 40));
5121 else if (sctx->gfx_level == GFX9)
5122 radeon_set_uconfig_reg(R_030944_VGT_TF_MEMORY_BASE_HI, S_030944_BASE_HI(factor_va >> 40));
5123 } else {
5124 radeon_set_config_reg(R_008988_VGT_TF_RING_SIZE, S_008988_SIZE(tf_ring_size_field));
5125 radeon_set_config_reg(R_0089B8_VGT_TF_MEMORY_BASE, factor_va >> 8);
5126 radeon_set_config_reg(R_0089B0_VGT_HS_OFFCHIP_PARAM, sscreen->hs.hs_offchip_param);
5127 }
5128 radeon_end();
5129 }
5130
5131 if (sctx->gfx_level >= GFX11) {
5132 /* We must wait for idle using an EOP event before changing the attribute ring registers.
5133 * Use the bottom-of-pipe EOP event, but use the PWS TS counter instead of the counter
5134 * in memory.
5135 */
5136 si_cp_release_acquire_mem_pws(sctx, &sctx->gfx_cs, V_028A90_BOTTOM_OF_PIPE_TS, 0,
5137 V_580_CP_ME, 0);
5138
5139 uint64_t attr_address = sscreen->attribute_pos_prim_ring->gpu_address;
5140 assert((attr_address >> 32) == sscreen->info.address32_hi);
5141
5142 radeon_begin(&sctx->gfx_cs);
5143 radeon_set_uconfig_reg_seq(R_031110_SPI_GS_THROTTLE_CNTL1, 4);
5144 radeon_emit(0x12355123); /* SPI_GS_THROTTLE_CNTL1 */
5145 radeon_emit(0x1544D); /* SPI_GS_THROTTLE_CNTL2 */
5146 radeon_emit(attr_address >> 16); /* SPI_ATTRIBUTE_RING_BASE */
5147 radeon_emit(S_03111C_MEM_SIZE((sscreen->info.attribute_ring_size_per_se >> 16) - 1) |
5148 S_03111C_BIG_PAGE(sscreen->info.discardable_allows_big_page) |
5149 S_03111C_L1_POLICY(1)); /* SPI_ATTRIBUTE_RING_SIZE */
5150
5151 if (sctx->gfx_level >= GFX12) {
5152 uint64_t pos_address = attr_address + sscreen->info.pos_ring_offset;
5153 uint64_t prim_address = attr_address + sscreen->info.prim_ring_offset;
5154
5155 /* When one of these 4 registers is updated, all 4 must be updated. */
5156 radeon_set_uconfig_reg_seq(R_0309A0_GE_POS_RING_BASE, 4);
5157 radeon_emit(pos_address >> 16); /* R_0309A0_GE_POS_RING_BASE */
5158 radeon_emit(S_0309A4_MEM_SIZE(sscreen->info.pos_ring_size_per_se >> 5)); /* R_0309A4_GE_POS_RING_SIZE */
5159 radeon_emit(prim_address >> 16); /* R_0309A8_GE_PRIM_RING_BASE */
5160 radeon_emit(S_0309AC_MEM_SIZE(sscreen->info.prim_ring_size_per_se >> 5) |
5161 S_0309AC_SCOPE(gfx12_scope_device) |
5162 S_0309AC_PAF_TEMPORAL(gfx12_store_high_temporal_stay_dirty) |
5163 S_0309AC_PAB_TEMPORAL(gfx12_load_last_use_discard) |
5164 S_0309AC_SPEC_DATA_READ(gfx12_spec_read_auto) |
5165 S_0309AC_FORCE_SE_SCOPE(1) |
5166 S_0309AC_PAB_NOFILL(1)); /* R_0309AC_GE_PRIM_RING_SIZE */
5167 }
5168 radeon_end();
5169 }
5170 }
5171
si_init_shader_functions(struct si_context * sctx)5172 void si_init_shader_functions(struct si_context *sctx)
5173 {
5174 sctx->atoms.s.vgt_pipeline_state.emit = si_emit_vgt_pipeline_state;
5175 sctx->atoms.s.scratch_state.emit = si_emit_scratch_state;
5176 sctx->atoms.s.spi_ge_ring_state.emit = si_emit_spi_ge_ring_state;
5177
5178 if (sctx->gfx_level >= GFX12)
5179 sctx->atoms.s.tess_io_layout.emit = gfx12_emit_tess_io_layout_state;
5180 else
5181 sctx->atoms.s.tess_io_layout.emit = gfx6_emit_tess_io_layout_state;
5182
5183 sctx->b.create_vs_state = si_create_shader;
5184 sctx->b.create_tcs_state = si_create_shader;
5185 sctx->b.create_tes_state = si_create_shader;
5186 sctx->b.create_gs_state = si_create_shader;
5187 sctx->b.create_fs_state = si_create_shader;
5188
5189 sctx->b.bind_vs_state = si_bind_vs_shader;
5190 sctx->b.bind_tcs_state = si_bind_tcs_shader;
5191 sctx->b.bind_tes_state = si_bind_tes_shader;
5192 sctx->b.bind_gs_state = si_bind_gs_shader;
5193 sctx->b.bind_fs_state = si_bind_ps_shader;
5194
5195 sctx->b.delete_vs_state = si_delete_shader_selector;
5196 sctx->b.delete_tcs_state = si_delete_shader_selector;
5197 sctx->b.delete_tes_state = si_delete_shader_selector;
5198 sctx->b.delete_gs_state = si_delete_shader_selector;
5199 sctx->b.delete_fs_state = si_delete_shader_selector;
5200
5201 sctx->b.set_patch_vertices = si_set_patch_vertices;
5202
5203 /* This unrolls the loops in si_emit_spi_map and inlines memcmp and memcpys.
5204 * It improves performance for viewperf/snx.
5205 */
5206 sctx->emit_spi_map[0] = si_emit_spi_map<0>;
5207 sctx->emit_spi_map[1] = si_emit_spi_map<1>;
5208 sctx->emit_spi_map[2] = si_emit_spi_map<2>;
5209 sctx->emit_spi_map[3] = si_emit_spi_map<3>;
5210 sctx->emit_spi_map[4] = si_emit_spi_map<4>;
5211 sctx->emit_spi_map[5] = si_emit_spi_map<5>;
5212 sctx->emit_spi_map[6] = si_emit_spi_map<6>;
5213 sctx->emit_spi_map[7] = si_emit_spi_map<7>;
5214 sctx->emit_spi_map[8] = si_emit_spi_map<8>;
5215 sctx->emit_spi_map[9] = si_emit_spi_map<9>;
5216 sctx->emit_spi_map[10] = si_emit_spi_map<10>;
5217 sctx->emit_spi_map[11] = si_emit_spi_map<11>;
5218 sctx->emit_spi_map[12] = si_emit_spi_map<12>;
5219 sctx->emit_spi_map[13] = si_emit_spi_map<13>;
5220 sctx->emit_spi_map[14] = si_emit_spi_map<14>;
5221 sctx->emit_spi_map[15] = si_emit_spi_map<15>;
5222 sctx->emit_spi_map[16] = si_emit_spi_map<16>;
5223 sctx->emit_spi_map[17] = si_emit_spi_map<17>;
5224 sctx->emit_spi_map[18] = si_emit_spi_map<18>;
5225 sctx->emit_spi_map[19] = si_emit_spi_map<19>;
5226 sctx->emit_spi_map[20] = si_emit_spi_map<20>;
5227 sctx->emit_spi_map[21] = si_emit_spi_map<21>;
5228 sctx->emit_spi_map[22] = si_emit_spi_map<22>;
5229 sctx->emit_spi_map[23] = si_emit_spi_map<23>;
5230 sctx->emit_spi_map[24] = si_emit_spi_map<24>;
5231 sctx->emit_spi_map[25] = si_emit_spi_map<25>;
5232 sctx->emit_spi_map[26] = si_emit_spi_map<26>;
5233 sctx->emit_spi_map[27] = si_emit_spi_map<27>;
5234 sctx->emit_spi_map[28] = si_emit_spi_map<28>;
5235 sctx->emit_spi_map[29] = si_emit_spi_map<29>;
5236 sctx->emit_spi_map[30] = si_emit_spi_map<30>;
5237 sctx->emit_spi_map[31] = si_emit_spi_map<31>;
5238 sctx->emit_spi_map[32] = si_emit_spi_map<32>;
5239 }
5240