1 /*
2  * Copyright © 2016 Red Hat.
3  * Copyright © 2016 Bas Nieuwenhuizen
4  *
5  * based in part on anv driver which is:
6  * Copyright © 2015 Intel Corporation
7  *
8  * Permission is hereby granted, free of charge, to any person obtaining a
9  * copy of this software and associated documentation files (the "Software"),
10  * to deal in the Software without restriction, including without limitation
11  * the rights to use, copy, modify, merge, publish, distribute, sublicense,
12  * and/or sell copies of the Software, and to permit persons to whom the
13  * Software is furnished to do so, subject to the following conditions:
14  *
15  * The above copyright notice and this permission notice (including the next
16  * paragraph) shall be included in all copies or substantial portions of the
17  * Software.
18  *
19  * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
20  * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
21  * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL
22  * THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
23  * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING
24  * FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS
25  * IN THE SOFTWARE.
26  */
27 
28 #ifndef RADV_SHADER_H
29 #define RADV_SHADER_H
30 
31 #include "ac_binary.h"
32 #include "ac_shader_util.h"
33 
34 #include "amd_family.h"
35 #include "radv_constants.h"
36 
37 #include "nir/nir.h"
38 #include "vulkan/util/vk_object.h"
39 #include "vulkan/util/vk_shader_module.h"
40 #include "vulkan/vulkan.h"
41 
42 #define RADV_VERT_ATTRIB_MAX MAX2(VERT_ATTRIB_MAX, VERT_ATTRIB_GENERIC0 + MAX_VERTEX_ATTRIBS)
43 
44 struct radv_physical_device;
45 struct radv_device;
46 struct radv_pipeline;
47 struct radv_pipeline_cache;
48 struct radv_pipeline_key;
49 struct radv_vs_input_state;
50 
51 enum radv_vs_input_alpha_adjust {
52    ALPHA_ADJUST_NONE = 0,
53    ALPHA_ADJUST_SNORM = 1,
54    ALPHA_ADJUST_SSCALED = 2,
55    ALPHA_ADJUST_SINT = 3,
56 };
57 
58 struct radv_pipeline_key {
59    uint32_t has_multiview_view_index : 1;
60    uint32_t optimisations_disabled : 1;
61    uint32_t invariant_geom : 1;
62    uint32_t use_ngg : 1;
63 
64    struct {
65       uint32_t instance_rate_inputs;
66       uint32_t instance_rate_divisors[MAX_VERTEX_ATTRIBS];
67       uint8_t vertex_attribute_formats[MAX_VERTEX_ATTRIBS];
68       uint32_t vertex_attribute_bindings[MAX_VERTEX_ATTRIBS];
69       uint32_t vertex_attribute_offsets[MAX_VERTEX_ATTRIBS];
70       uint32_t vertex_attribute_strides[MAX_VERTEX_ATTRIBS];
71       uint8_t vertex_binding_align[MAX_VBS];
72       enum radv_vs_input_alpha_adjust vertex_alpha_adjust[MAX_VERTEX_ATTRIBS];
73       uint32_t vertex_post_shuffle;
74       uint32_t provoking_vtx_last : 1;
75       uint32_t dynamic_input_state : 1;
76       uint8_t topology;
77    } vs;
78 
79    struct {
80       unsigned tess_input_vertices;
81    } tcs;
82 
83    struct {
84       uint32_t col_format;
85       uint32_t is_int8;
86       uint32_t is_int10;
87       uint8_t log2_ps_iter_samples;
88       uint8_t num_samples;
89 
90       bool lower_discard_to_demote;
91       bool enable_mrt_output_nan_fixup;
92       uint8_t force_vrs;
93    } ps;
94 
95    struct {
96       /* Non-zero if a required subgroup size is specified via
97        * VK_EXT_subgroup_size_control.
98        */
99       uint8_t compute_subgroup_size;
100       bool require_full_subgroups;
101    } cs;
102 };
103 
104 enum radv_compiler_debug_level {
105    RADV_COMPILER_DEBUG_LEVEL_PERFWARN,
106    RADV_COMPILER_DEBUG_LEVEL_ERROR,
107 };
108 
109 struct radv_nir_compiler_options {
110    struct radv_pipeline_layout *layout;
111    struct radv_pipeline_key key;
112    bool explicit_scratch_args;
113    bool robust_buffer_access;
114    bool adjust_frag_coord_z;
115    bool dump_shader;
116    bool dump_preoptir;
117    bool record_ir;
118    bool record_stats;
119    bool check_ir;
120    bool has_ls_vgpr_init_bug;
121    bool has_image_load_dcc_bug;
122    bool enable_mrt_output_nan_fixup;
123    bool wgp_mode;
124    bool remap_spi_ps_input;
125    enum radeon_family family;
126    enum chip_class chip_class;
127    const struct radeon_info *info;
128    uint32_t address32_hi;
129    uint8_t force_vrs_rates;
130 
131    struct {
132       void (*func)(void *private_data, enum radv_compiler_debug_level level, const char *message);
133       void *private_data;
134    } debug;
135 };
136 
137 enum radv_ud_index {
138    AC_UD_SCRATCH_RING_OFFSETS = 0,
139    AC_UD_PUSH_CONSTANTS = 1,
140    AC_UD_INLINE_PUSH_CONSTANTS = 2,
141    AC_UD_INDIRECT_DESCRIPTOR_SETS = 3,
142    AC_UD_VIEW_INDEX = 4,
143    AC_UD_STREAMOUT_BUFFERS = 5,
144    AC_UD_NGG_GS_STATE = 6,
145    AC_UD_NGG_CULLING_SETTINGS = 7,
146    AC_UD_NGG_VIEWPORT = 8,
147    AC_UD_SHADER_START = 9,
148    AC_UD_VS_VERTEX_BUFFERS = AC_UD_SHADER_START,
149    AC_UD_VS_BASE_VERTEX_START_INSTANCE,
150    AC_UD_VS_PROLOG_INPUTS,
151    AC_UD_VS_MAX_UD,
152    AC_UD_PS_MAX_UD,
153    AC_UD_CS_GRID_SIZE = AC_UD_SHADER_START,
154    AC_UD_CS_SBT_DESCRIPTORS,
155    AC_UD_CS_RAY_LAUNCH_SIZE,
156    AC_UD_CS_MAX_UD,
157    AC_UD_GS_MAX_UD,
158    AC_UD_TCS_MAX_UD,
159    AC_UD_TES_MAX_UD,
160    AC_UD_MAX_UD = AC_UD_TCS_MAX_UD,
161 };
162 
163 struct radv_stream_output {
164    uint8_t location;
165    uint8_t buffer;
166    uint16_t offset;
167    uint8_t component_mask;
168    uint8_t stream;
169 };
170 
171 struct radv_streamout_info {
172    uint16_t num_outputs;
173    struct radv_stream_output outputs[MAX_SO_OUTPUTS];
174    uint16_t strides[MAX_SO_BUFFERS];
175    uint32_t enabled_stream_buffers_mask;
176 };
177 
178 struct radv_userdata_info {
179    int8_t sgpr_idx;
180    uint8_t num_sgprs;
181 };
182 
183 struct radv_userdata_locations {
184    struct radv_userdata_info descriptor_sets[MAX_SETS];
185    struct radv_userdata_info shader_data[AC_UD_MAX_UD];
186    uint32_t descriptor_sets_enabled;
187 };
188 
189 struct radv_vs_output_info {
190    uint8_t vs_output_param_offset[VARYING_SLOT_MAX];
191    uint8_t clip_dist_mask;
192    uint8_t cull_dist_mask;
193    uint8_t param_exports;
194    bool writes_pointsize;
195    bool writes_layer;
196    bool writes_viewport_index;
197    bool writes_primitive_shading_rate;
198    bool export_prim_id;
199    bool export_clip_dists;
200    unsigned pos_exports;
201 };
202 
203 struct radv_es_output_info {
204    uint32_t esgs_itemsize;
205 };
206 
207 struct gfx9_gs_info {
208    uint32_t vgt_gs_onchip_cntl;
209    uint32_t vgt_gs_max_prims_per_subgroup;
210    uint32_t vgt_esgs_ring_itemsize;
211    uint32_t lds_size;
212 };
213 
214 struct gfx10_ngg_info {
215    uint16_t ngg_emit_size; /* in dwords */
216    uint32_t hw_max_esverts;
217    uint32_t max_gsprims;
218    uint32_t max_out_verts;
219    uint32_t prim_amp_factor;
220    uint32_t vgt_esgs_ring_itemsize;
221    uint32_t esgs_ring_size;
222    bool max_vert_out_per_gs_instance;
223    bool enable_vertex_grouping;
224 };
225 
226 struct radv_shader_info {
227    bool loads_push_constants;
228    bool loads_dynamic_offsets;
229    uint8_t min_push_constant_used;
230    uint8_t max_push_constant_used;
231    bool has_only_32bit_push_constants;
232    bool has_indirect_push_constants;
233    uint32_t desc_set_used_mask;
234    bool uses_view_index;
235    bool uses_invocation_id;
236    bool uses_prim_id;
237    uint8_t wave_size;
238    uint8_t ballot_bit_size;
239    struct radv_userdata_locations user_sgprs_locs;
240    unsigned num_user_sgprs;
241    unsigned num_input_sgprs;
242    unsigned num_input_vgprs;
243    bool is_ngg;
244    bool is_ngg_passthrough;
245    bool has_ngg_culling;
246    bool has_ngg_early_prim_export;
247    uint32_t num_lds_blocks_when_not_culling;
248    uint32_t num_tess_patches;
249    unsigned workgroup_size;
250    struct {
251       uint8_t input_usage_mask[RADV_VERT_ATTRIB_MAX];
252       uint8_t output_usage_mask[VARYING_SLOT_VAR31 + 1];
253       bool needs_draw_id;
254       bool needs_instance_id;
255       struct radv_vs_output_info outinfo;
256       struct radv_es_output_info es_info;
257       bool as_es;
258       bool as_ls;
259       bool tcs_in_out_eq;
260       uint64_t tcs_temp_only_input_mask;
261       uint8_t num_linked_outputs;
262       bool needs_base_instance;
263       bool use_per_attribute_vb_descs;
264       uint32_t vb_desc_usage_mask;
265       bool has_prolog;
266       bool dynamic_inputs;
267    } vs;
268    struct {
269       uint8_t output_usage_mask[VARYING_SLOT_VAR31 + 1];
270       uint8_t num_stream_output_components[4];
271       uint8_t output_streams[VARYING_SLOT_VAR31 + 1];
272       uint8_t max_stream;
273       unsigned gsvs_vertex_size;
274       unsigned max_gsvs_emit_size;
275       unsigned vertices_in;
276       unsigned vertices_out;
277       unsigned output_prim;
278       unsigned invocations;
279       unsigned es_type; /* GFX9: VS or TES */
280       uint8_t num_linked_inputs;
281    } gs;
282    struct {
283       uint8_t output_usage_mask[VARYING_SLOT_VAR31 + 1];
284       struct radv_vs_output_info outinfo;
285       struct radv_es_output_info es_info;
286       bool as_es;
287       unsigned primitive_mode;
288       enum gl_tess_spacing spacing;
289       bool ccw;
290       bool point_mode;
291       uint8_t num_linked_inputs;
292       uint8_t num_linked_patch_inputs;
293       uint8_t num_linked_outputs;
294    } tes;
295    struct {
296       bool uses_sample_shading;
297       bool needs_sample_positions;
298       bool writes_memory;
299       bool writes_z;
300       bool writes_stencil;
301       bool writes_sample_mask;
302       bool has_pcoord;
303       bool prim_id_input;
304       bool layer_input;
305       bool viewport_index_input;
306       uint8_t num_input_clips_culls;
307       uint32_t input_mask;
308       uint32_t flat_shaded_mask;
309       uint32_t explicit_shaded_mask;
310       uint32_t float16_shaded_mask;
311       uint32_t num_interp;
312       bool can_discard;
313       bool early_fragment_test;
314       bool post_depth_coverage;
315       bool reads_sample_mask_in;
316       bool reads_front_face;
317       bool reads_sample_id;
318       bool reads_frag_shading_rate;
319       bool reads_barycentric_model;
320       bool reads_persp_sample;
321       bool reads_persp_center;
322       bool reads_persp_centroid;
323       bool reads_linear_sample;
324       bool reads_linear_center;
325       bool reads_linear_centroid;
326       uint8_t reads_frag_coord_mask;
327       uint8_t reads_sample_pos_mask;
328       uint8_t depth_layout;
329       bool allow_flat_shading;
330       unsigned spi_ps_input;
331    } ps;
332    struct {
333       bool uses_grid_size;
334       bool uses_block_id[3];
335       bool uses_thread_id[3];
336       bool uses_local_invocation_idx;
337       unsigned block_size[3];
338 
339       uint8_t subgroup_size;
340 
341       bool uses_sbt;
342       bool uses_ray_launch_size;
343    } cs;
344    struct {
345       uint64_t tes_inputs_read;
346       uint64_t tes_patch_inputs_read;
347       unsigned tcs_vertices_out;
348       uint32_t num_lds_blocks;
349       uint8_t num_linked_inputs;
350       uint8_t num_linked_outputs;
351       uint8_t num_linked_patch_outputs;
352       bool tes_reads_tess_factors : 1;
353    } tcs;
354 
355    struct radv_streamout_info so;
356 
357    struct gfx9_gs_info gs_ring_info;
358    struct gfx10_ngg_info ngg_info;
359 };
360 
361 struct radv_vs_input_state {
362    uint32_t attribute_mask;
363    uint32_t misaligned_mask;
364    uint32_t possibly_misaligned_mask;
365 
366    uint32_t instance_rate_inputs;
367    uint32_t nontrivial_divisors;
368    uint32_t post_shuffle;
369    /* Having two separate fields instead of a single uint64_t makes it easier to remove attributes
370     * using bitwise arithmetic.
371     */
372    uint32_t alpha_adjust_lo;
373    uint32_t alpha_adjust_hi;
374 
375    uint8_t bindings[MAX_VERTEX_ATTRIBS];
376    uint32_t divisors[MAX_VERTEX_ATTRIBS];
377    uint32_t offsets[MAX_VERTEX_ATTRIBS];
378    uint8_t formats[MAX_VERTEX_ATTRIBS];
379    uint8_t format_align_req_minus_1[MAX_VERTEX_ATTRIBS];
380    uint8_t format_sizes[MAX_VERTEX_ATTRIBS];
381 };
382 
383 struct radv_vs_prolog_key {
384    const struct radv_vs_input_state *state;
385    unsigned num_attributes;
386    uint32_t misaligned_mask;
387    bool as_ls;
388    bool is_ngg;
389    bool wave32;
390    gl_shader_stage next_stage;
391 };
392 
393 enum radv_shader_binary_type { RADV_BINARY_TYPE_LEGACY, RADV_BINARY_TYPE_RTLD };
394 
395 struct radv_shader_binary {
396    enum radv_shader_binary_type type;
397    gl_shader_stage stage;
398    bool is_gs_copy_shader;
399 
400    struct ac_shader_config config;
401    struct radv_shader_info info;
402 
403    /* Self-referential size so we avoid consistency issues. */
404    uint32_t total_size;
405 };
406 
407 struct radv_shader_binary_legacy {
408    struct radv_shader_binary base;
409    unsigned code_size;
410    unsigned exec_size;
411    unsigned ir_size;
412    unsigned disasm_size;
413    unsigned stats_size;
414 
415    /* data has size of stats_size + code_size + ir_size + disasm_size + 2,
416     * where the +2 is for 0 of the ir strings. */
417    uint8_t data[0];
418 };
419 
420 struct radv_shader_binary_rtld {
421    struct radv_shader_binary base;
422    unsigned elf_size;
423    unsigned llvm_ir_size;
424    uint8_t data[0];
425 };
426 
427 struct radv_prolog_binary {
428    uint8_t num_sgprs;
429    uint8_t num_vgprs;
430    uint8_t num_preserved_sgprs;
431    unsigned code_size;
432    uint8_t data[0];
433 };
434 
435 struct radv_shader_arena {
436    struct list_head list;
437    struct list_head entries;
438    struct radeon_winsys_bo *bo;
439    char *ptr;
440 };
441 
442 union radv_shader_arena_block {
443    struct list_head pool;
444    struct {
445       /* List of blocks in the arena, sorted by address. */
446       struct list_head list;
447       /* For holes, a list_head for the free-list. For allocations, freelist.prev=NULL and
448        * freelist.next is a pointer associated with the allocation.
449        */
450       struct list_head freelist;
451       struct radv_shader_arena *arena;
452       uint32_t offset;
453       uint32_t size;
454    };
455 };
456 
457 struct radv_shader_variant {
458    uint32_t ref_count;
459 
460    struct radeon_winsys_bo *bo;
461    union radv_shader_arena_block *alloc;
462    struct ac_shader_config config;
463    uint8_t *code_ptr;
464    uint32_t code_size;
465    uint32_t exec_size;
466    struct radv_shader_info info;
467 
468    /* debug only */
469    char *spirv;
470    uint32_t spirv_size;
471    char *nir_string;
472    char *disasm_string;
473    char *ir_string;
474    uint32_t *statistics;
475 };
476 
477 struct radv_shader_prolog {
478    struct radeon_winsys_bo *bo;
479    union radv_shader_arena_block *alloc;
480    uint32_t rsrc1;
481    uint8_t num_preserved_sgprs;
482    bool nontrivial_divisors;
483 };
484 
485 void radv_optimize_nir(const struct radv_device *device, struct nir_shader *shader,
486                        bool optimize_conservatively, bool allow_copies);
487 void radv_optimize_nir_algebraic(nir_shader *shader, bool opt_offsets);
488 bool radv_nir_lower_ycbcr_textures(nir_shader *shader, const struct radv_pipeline_layout *layout);
489 
490 nir_shader *radv_shader_compile_to_nir(struct radv_device *device, struct vk_shader_module *module,
491                                        const char *entrypoint_name, gl_shader_stage stage,
492                                        const VkSpecializationInfo *spec_info,
493                                        const struct radv_pipeline_layout *layout,
494                                        const struct radv_pipeline_key *key);
495 
496 void radv_init_shader_arenas(struct radv_device *device);
497 void radv_destroy_shader_arenas(struct radv_device *device);
498 
499 VkResult radv_create_shaders(struct radv_pipeline *pipeline,
500                              struct radv_pipeline_layout *pipeline_layout,
501                              struct radv_device *device, struct radv_pipeline_cache *cache,
502                              const struct radv_pipeline_key *key,
503                              const VkPipelineShaderStageCreateInfo **pStages,
504                              const VkPipelineCreateFlags flags, const uint8_t *custom_hash,
505                              VkPipelineCreationFeedbackEXT *pipeline_feedback,
506                              VkPipelineCreationFeedbackEXT **stage_feedbacks);
507 
508 struct radv_shader_variant *radv_shader_variant_create(struct radv_device *device,
509                                                        const struct radv_shader_binary *binary,
510                                                        bool keep_shader_info, bool from_cache);
511 struct radv_shader_variant *radv_shader_variant_compile(
512    struct radv_device *device, struct vk_shader_module *module, struct nir_shader *const *shaders,
513    int shader_count, struct radv_pipeline_layout *layout, const struct radv_pipeline_key *key,
514    struct radv_shader_info *info, bool keep_shader_info, bool keep_statistic_info,
515    struct radv_shader_binary **binary_out);
516 
517 struct radv_shader_variant *
518 radv_create_gs_copy_shader(struct radv_device *device, struct nir_shader *nir,
519                            struct radv_shader_info *info, struct radv_shader_binary **binary_out,
520                            bool multiview, bool keep_shader_info, bool keep_statistic_info,
521                            bool disable_optimizations);
522 
523 struct radv_shader_variant *radv_create_trap_handler_shader(struct radv_device *device);
524 
525 struct radv_shader_prolog *radv_create_vs_prolog(struct radv_device *device,
526                                                  const struct radv_vs_prolog_key *key);
527 
528 void radv_shader_variant_destroy(struct radv_device *device, struct radv_shader_variant *variant);
529 
530 void radv_prolog_destroy(struct radv_device *device, struct radv_shader_prolog *prolog);
531 
532 uint64_t radv_shader_variant_get_va(const struct radv_shader_variant *variant);
533 struct radv_shader_variant *radv_find_shader_variant(struct radv_device *device, uint64_t pc);
534 
535 unsigned radv_get_max_waves(const struct radv_device *device, struct radv_shader_variant *variant,
536                             gl_shader_stage stage);
537 
538 const char *radv_get_shader_name(struct radv_shader_info *info, gl_shader_stage stage);
539 
540 unsigned radv_compute_spi_ps_input(const struct radv_device *device,
541                                    const struct radv_shader_info *info);
542 
543 bool radv_can_dump_shader(struct radv_device *device, struct vk_shader_module *module,
544                           bool meta_shader);
545 
546 bool radv_can_dump_shader_stats(struct radv_device *device, struct vk_shader_module *module);
547 
548 VkResult radv_dump_shader_stats(struct radv_device *device, struct radv_pipeline *pipeline,
549                                 gl_shader_stage stage, FILE *output);
550 
551 static inline unsigned
calculate_tess_lds_size(enum chip_class chip_class,unsigned tcs_num_input_vertices,unsigned tcs_num_output_vertices,unsigned tcs_num_inputs,unsigned tcs_num_patches,unsigned tcs_num_outputs,unsigned tcs_num_patch_outputs)552 calculate_tess_lds_size(enum chip_class chip_class, unsigned tcs_num_input_vertices,
553                         unsigned tcs_num_output_vertices, unsigned tcs_num_inputs,
554                         unsigned tcs_num_patches, unsigned tcs_num_outputs,
555                         unsigned tcs_num_patch_outputs)
556 {
557    unsigned input_vertex_size = tcs_num_inputs * 16;
558    unsigned output_vertex_size = tcs_num_outputs * 16;
559 
560    unsigned input_patch_size = tcs_num_input_vertices * input_vertex_size;
561 
562    unsigned pervertex_output_patch_size = tcs_num_output_vertices * output_vertex_size;
563    unsigned output_patch_size = pervertex_output_patch_size + tcs_num_patch_outputs * 16;
564 
565    unsigned output_patch0_offset = input_patch_size * tcs_num_patches;
566 
567    unsigned lds_size = output_patch0_offset + output_patch_size * tcs_num_patches;
568 
569    if (chip_class >= GFX7) {
570       assert(lds_size <= 65536);
571       lds_size = align(lds_size, 512) / 512;
572    } else {
573       assert(lds_size <= 32768);
574       lds_size = align(lds_size, 256) / 256;
575    }
576 
577    return lds_size;
578 }
579 
580 static inline unsigned
get_tcs_num_patches(unsigned tcs_num_input_vertices,unsigned tcs_num_output_vertices,unsigned tcs_num_inputs,unsigned tcs_num_outputs,unsigned tcs_num_patch_outputs,unsigned tess_offchip_block_dw_size,enum chip_class chip_class,enum radeon_family family)581 get_tcs_num_patches(unsigned tcs_num_input_vertices, unsigned tcs_num_output_vertices,
582                     unsigned tcs_num_inputs, unsigned tcs_num_outputs,
583                     unsigned tcs_num_patch_outputs, unsigned tess_offchip_block_dw_size,
584                     enum chip_class chip_class, enum radeon_family family)
585 {
586    uint32_t input_vertex_size = tcs_num_inputs * 16;
587    uint32_t input_patch_size = tcs_num_input_vertices * input_vertex_size;
588    uint32_t output_vertex_size = tcs_num_outputs * 16;
589    uint32_t pervertex_output_patch_size = tcs_num_output_vertices * output_vertex_size;
590    uint32_t output_patch_size = pervertex_output_patch_size + tcs_num_patch_outputs * 16;
591 
592    /* Ensure that we only need one wave per SIMD so we don't need to check
593     * resource usage. Also ensures that the number of tcs in and out
594     * vertices per threadgroup are at most 256.
595     */
596    unsigned num_patches = 64 / MAX2(tcs_num_input_vertices, tcs_num_output_vertices) * 4;
597    /* Make sure that the data fits in LDS. This assumes the shaders only
598     * use LDS for the inputs and outputs.
599     */
600    unsigned hardware_lds_size = 32768;
601 
602    /* Looks like STONEY hangs if we use more than 32 KiB LDS in a single
603     * threadgroup, even though there is more than 32 KiB LDS.
604     *
605     * Test: dEQP-VK.tessellation.shader_input_output.barrier
606     */
607    if (chip_class >= GFX7 && family != CHIP_STONEY)
608       hardware_lds_size = 65536;
609 
610    if (input_patch_size + output_patch_size)
611       num_patches = MIN2(num_patches, hardware_lds_size / (input_patch_size + output_patch_size));
612    /* Make sure the output data fits in the offchip buffer */
613    if (output_patch_size)
614       num_patches = MIN2(num_patches, (tess_offchip_block_dw_size * 4) / output_patch_size);
615    /* Not necessary for correctness, but improves performance. The
616     * specific value is taken from the proprietary driver.
617     */
618    num_patches = MIN2(num_patches, 40);
619 
620    /* GFX6 bug workaround - limit LS-HS threadgroups to only one wave. */
621    if (chip_class == GFX6) {
622       unsigned one_wave = 64 / MAX2(tcs_num_input_vertices, tcs_num_output_vertices);
623       num_patches = MIN2(num_patches, one_wave);
624    }
625    return num_patches;
626 }
627 
628 void radv_lower_io(struct radv_device *device, nir_shader *nir);
629 
630 bool radv_lower_io_to_mem(struct radv_device *device, struct nir_shader *nir,
631                           const struct radv_shader_info *info, const struct radv_pipeline_key *pl_key);
632 
633 void radv_lower_ngg(struct radv_device *device, struct nir_shader *nir,
634                     const struct radv_shader_info *info,
635                     const struct radv_pipeline_key *pl_key);
636 
637 bool radv_consider_culling(struct radv_device *device, struct nir_shader *nir,
638                            uint64_t ps_inputs_read, unsigned num_vertices_per_primitive,
639                            const struct radv_shader_info *info);
640 
641 void radv_get_nir_options(struct radv_physical_device *device);
642 
643 #endif
644