projects / mesa.git / blob
? search:
summary | shortlog | log | commit | commitdiff | tree
history | raw | HEAD
radv: rework how the number of VGPRs is computed
[mesa.git] / src / amd / vulkan / radv_shader.c
1 /*
2 * Copyright © 2016 Red Hat.
3 * Copyright © 2016 Bas Nieuwenhuizen
4 *
5 * based in part on anv driver which is:
6 * Copyright © 2015 Intel Corporation
7 *
8 * Permission is hereby granted, free of charge, to any person obtaining a
9 * copy of this software and associated documentation files (the "Software"),
10 * to deal in the Software without restriction, including without limitation
11 * the rights to use, copy, modify, merge, publish, distribute, sublicense,
12 * and/or sell copies of the Software, and to permit persons to whom the
13 * Software is furnished to do so, subject to the following conditions:
14 *
15 * The above copyright notice and this permission notice (including the next
16 * paragraph) shall be included in all copies or substantial portions of the
17 * Software.
18 *
19 * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
20 * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
21 * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL
22 * THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
23 * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING
24 * FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS
25 * IN THE SOFTWARE.
26 */
27
28 #include "util/mesa-sha1.h"
29 #include "util/u_atomic.h"
30 #include "radv_debug.h"
31 #include "radv_private.h"
32 #include "radv_shader.h"
33 #include "radv_shader_helper.h"
34 #include "nir/nir.h"
35 #include "nir/nir_builder.h"
36 #include "spirv/nir_spirv.h"
37
38 #include <llvm-c/Core.h>
39 #include <llvm-c/TargetMachine.h>
40 #include <llvm-c/Support.h>
41
42 #include "sid.h"
43 #include "ac_binary.h"
44 #include "ac_llvm_util.h"
45 #include "ac_nir_to_llvm.h"
46 #include "vk_format.h"
47 #include "util/debug.h"
48 #include "ac_exp_param.h"
49
50 #include "util/string_buffer.h"
51
52 static const struct nir_shader_compiler_options nir_options = {
53 .vertex_id_zero_based = true,
54 .lower_scmp = true,
55 .lower_flrp16 = true,
56 .lower_flrp32 = true,
57 .lower_flrp64 = true,
58 .lower_device_index_to_zero = true,
59 .lower_fsat = true,
60 .lower_fdiv = true,
61 .lower_bitfield_insert_to_bitfield_select = true,
62 .lower_bitfield_extract = true,
63 .lower_sub = true,
64 .lower_pack_snorm_2x16 = true,
65 .lower_pack_snorm_4x8 = true,
66 .lower_pack_unorm_2x16 = true,
67 .lower_pack_unorm_4x8 = true,
68 .lower_unpack_snorm_2x16 = true,
69 .lower_unpack_snorm_4x8 = true,
70 .lower_unpack_unorm_2x16 = true,
71 .lower_unpack_unorm_4x8 = true,
72 .lower_extract_byte = true,
73 .lower_extract_word = true,
74 .lower_ffma = true,
75 .lower_fpow = true,
76 .lower_mul_2x32_64 = true,
77 .max_unroll_iterations = 32
78 };
79
80 VkResult radv_CreateShaderModule(
81 VkDevice _device,
82 const VkShaderModuleCreateInfo* pCreateInfo,
83 const VkAllocationCallbacks* pAllocator,
84 VkShaderModule* pShaderModule)
85 {
86 RADV_FROM_HANDLE(radv_device, device, _device);
87 struct radv_shader_module *module;
88
89 assert(pCreateInfo->sType == VK_STRUCTURE_TYPE_SHADER_MODULE_CREATE_INFO);
90 assert(pCreateInfo->flags == 0);
91
92 module = vk_alloc2(&device->alloc, pAllocator,
93 sizeof(*module) + pCreateInfo->codeSize, 8,
94 VK_SYSTEM_ALLOCATION_SCOPE_OBJECT);
95 if (module == NULL)
96 return vk_error(device->instance, VK_ERROR_OUT_OF_HOST_MEMORY);
97
98 module->nir = NULL;
99 module->size = pCreateInfo->codeSize;
100 memcpy(module->data, pCreateInfo->pCode, module->size);
101
102 _mesa_sha1_compute(module->data, module->size, module->sha1);
103
104 *pShaderModule = radv_shader_module_to_handle(module);
105
106 return VK_SUCCESS;
107 }
108
109 void radv_DestroyShaderModule(
110 VkDevice _device,
111 VkShaderModule _module,
112 const VkAllocationCallbacks* pAllocator)
113 {
114 RADV_FROM_HANDLE(radv_device, device, _device);
115 RADV_FROM_HANDLE(radv_shader_module, module, _module);
116
117 if (!module)
118 return;
119
120 vk_free2(&device->alloc, pAllocator, module);
121 }
122
123 void
124 radv_optimize_nir(struct nir_shader *shader, bool optimize_conservatively,
125 bool allow_copies)
126 {
127 bool progress;
128 unsigned lower_flrp =
129 (shader->options->lower_flrp16 ? 16 : 0) |
130 (shader->options->lower_flrp32 ? 32 : 0) |
131 (shader->options->lower_flrp64 ? 64 : 0);
132
133 do {
134 progress = false;
135
136 NIR_PASS(progress, shader, nir_split_array_vars, nir_var_function_temp);
137 NIR_PASS(progress, shader, nir_shrink_vec_array_vars, nir_var_function_temp);
138
139 NIR_PASS_V(shader, nir_lower_vars_to_ssa);
140 NIR_PASS_V(shader, nir_lower_pack);
141
142 if (allow_copies) {
143 /* Only run this pass in the first call to
144 * radv_optimize_nir. Later calls assume that we've
145 * lowered away any copy_deref instructions and we
146 * don't want to introduce any more.
147 */
148 NIR_PASS(progress, shader, nir_opt_find_array_copies);
149 }
150
151 NIR_PASS(progress, shader, nir_opt_copy_prop_vars);
152 NIR_PASS(progress, shader, nir_opt_dead_write_vars);
153
154 NIR_PASS_V(shader, nir_lower_alu_to_scalar, NULL);
155 NIR_PASS_V(shader, nir_lower_phis_to_scalar);
156
157 NIR_PASS(progress, shader, nir_copy_prop);
158 NIR_PASS(progress, shader, nir_opt_remove_phis);
159 NIR_PASS(progress, shader, nir_opt_dce);
160 if (nir_opt_trivial_continues(shader)) {
161 progress = true;
162 NIR_PASS(progress, shader, nir_copy_prop);
163 NIR_PASS(progress, shader, nir_opt_remove_phis);
164 NIR_PASS(progress, shader, nir_opt_dce);
165 }
166 NIR_PASS(progress, shader, nir_opt_if, true);
167 NIR_PASS(progress, shader, nir_opt_dead_cf);
168 NIR_PASS(progress, shader, nir_opt_cse);
169 NIR_PASS(progress, shader, nir_opt_peephole_select, 8, true, true);
170 NIR_PASS(progress, shader, nir_opt_constant_folding);
171 NIR_PASS(progress, shader, nir_opt_algebraic);
172
173 if (lower_flrp != 0) {
174 bool lower_flrp_progress = false;
175 NIR_PASS(lower_flrp_progress,
176 shader,
177 nir_lower_flrp,
178 lower_flrp,
179 false /* always_precise */,
180 shader->options->lower_ffma);
181 if (lower_flrp_progress) {
182 NIR_PASS(progress, shader,
183 nir_opt_constant_folding);
184 progress = true;
185 }
186
187 /* Nothing should rematerialize any flrps, so we only
188 * need to do this lowering once.
189 */
190 lower_flrp = 0;
191 }
192
193 NIR_PASS(progress, shader, nir_opt_undef);
194 NIR_PASS(progress, shader, nir_opt_conditional_discard);
195 if (shader->options->max_unroll_iterations) {
196 NIR_PASS(progress, shader, nir_opt_loop_unroll, 0);
197 }
198 } while (progress && !optimize_conservatively);
199
200 NIR_PASS(progress, shader, nir_opt_shrink_load);
201 NIR_PASS(progress, shader, nir_opt_move_load_ubo);
202 }
203
204 nir_shader *
205 radv_shader_compile_to_nir(struct radv_device *device,
206 struct radv_shader_module *module,
207 const char *entrypoint_name,
208 gl_shader_stage stage,
209 const VkSpecializationInfo *spec_info,
210 const VkPipelineCreateFlags flags,
211 const struct radv_pipeline_layout *layout)
212 {
213 nir_shader *nir;
214 if (module->nir) {
215 /* Some things such as our meta clear/blit code will give us a NIR
216 * shader directly. In that case, we just ignore the SPIR-V entirely
217 * and just use the NIR shader */
218 nir = module->nir;
219 nir->options = &nir_options;
220 nir_validate_shader(nir, "in internal shader");
221
222 assert(exec_list_length(&nir->functions) == 1);
223 } else {
224 uint32_t *spirv = (uint32_t *) module->data;
225 assert(module->size % 4 == 0);
226
227 if (device->instance->debug_flags & RADV_DEBUG_DUMP_SPIRV)
228 radv_print_spirv(spirv, module->size, stderr);
229
230 uint32_t num_spec_entries = 0;
231 struct nir_spirv_specialization *spec_entries = NULL;
232 if (spec_info && spec_info->mapEntryCount > 0) {
233 num_spec_entries = spec_info->mapEntryCount;
234 spec_entries = malloc(num_spec_entries * sizeof(*spec_entries));
235 for (uint32_t i = 0; i < num_spec_entries; i++) {
236 VkSpecializationMapEntry entry = spec_info->pMapEntries[i];
237 const void *data = spec_info->pData + entry.offset;
238 assert(data + entry.size <= spec_info->pData + spec_info->dataSize);
239
240 spec_entries[i].id = spec_info->pMapEntries[i].constantID;
241 if (spec_info->dataSize == 8)
242 spec_entries[i].data64 = *(const uint64_t *)data;
243 else
244 spec_entries[i].data32 = *(const uint32_t *)data;
245 }
246 }
247 const struct spirv_to_nir_options spirv_options = {
248 .lower_ubo_ssbo_access_to_offsets = true,
249 .caps = {
250 .amd_gcn_shader = true,
251 .amd_shader_ballot = device->instance->perftest_flags & RADV_PERFTEST_SHADER_BALLOT,
252 .amd_trinary_minmax = true,
253 .derivative_group = true,
254 .descriptor_array_dynamic_indexing = true,
255 .descriptor_array_non_uniform_indexing = true,
256 .descriptor_indexing = true,
257 .device_group = true,
258 .draw_parameters = true,
259 .float16 = true,
260 .float64 = true,
261 .geometry_streams = true,
262 .image_read_without_format = true,
263 .image_write_without_format = true,
264 .int8 = true,
265 .int16 = true,
266 .int64 = true,
267 .int64_atomics = true,
268 .multiview = true,
269 .physical_storage_buffer_address = true,
270 .runtime_descriptor_array = true,
271 .shader_viewport_index_layer = true,
272 .stencil_export = true,
273 .storage_8bit = true,
274 .storage_16bit = true,
275 .storage_image_ms = true,
276 .subgroup_arithmetic = true,
277 .subgroup_ballot = true,
278 .subgroup_basic = true,
279 .subgroup_quad = true,
280 .subgroup_shuffle = true,
281 .subgroup_vote = true,
282 .tessellation = true,
283 .transform_feedback = true,
284 .variable_pointers = true,
285 },
286 .ubo_addr_format = nir_address_format_32bit_index_offset,
287 .ssbo_addr_format = nir_address_format_32bit_index_offset,
288 .phys_ssbo_addr_format = nir_address_format_64bit_global,
289 .push_const_addr_format = nir_address_format_logical,
290 .shared_addr_format = nir_address_format_32bit_offset,
291 };
292 nir = spirv_to_nir(spirv, module->size / 4,
293 spec_entries, num_spec_entries,
294 stage, entrypoint_name,
295 &spirv_options, &nir_options);
296 assert(nir->info.stage == stage);
297 nir_validate_shader(nir, "after spirv_to_nir");
298
299 free(spec_entries);
300
301 /* We have to lower away local constant initializers right before we
302 * inline functions. That way they get properly initialized at the top
303 * of the function and not at the top of its caller.
304 */
305 NIR_PASS_V(nir, nir_lower_constant_initializers, nir_var_function_temp);
306 NIR_PASS_V(nir, nir_lower_returns);
307 NIR_PASS_V(nir, nir_inline_functions);
308 NIR_PASS_V(nir, nir_opt_deref);
309
310 /* Pick off the single entrypoint that we want */
311 foreach_list_typed_safe(nir_function, func, node, &nir->functions) {
312 if (func->is_entrypoint)
313 func->name = ralloc_strdup(func, "main");
314 else
315 exec_node_remove(&func->node);
316 }
317 assert(exec_list_length(&nir->functions) == 1);
318
319 /* Make sure we lower constant initializers on output variables so that
320 * nir_remove_dead_variables below sees the corresponding stores
321 */
322 NIR_PASS_V(nir, nir_lower_constant_initializers, nir_var_shader_out);
323
324 /* Now that we've deleted all but the main function, we can go ahead and
325 * lower the rest of the constant initializers.
326 */
327 NIR_PASS_V(nir, nir_lower_constant_initializers, ~0);
328
329 /* Split member structs. We do this before lower_io_to_temporaries so that
330 * it doesn't lower system values to temporaries by accident.
331 */
332 NIR_PASS_V(nir, nir_split_var_copies);
333 NIR_PASS_V(nir, nir_split_per_member_structs);
334
335 NIR_PASS_V(nir, nir_remove_dead_variables,
336 nir_var_shader_in | nir_var_shader_out | nir_var_system_value);
337
338 NIR_PASS_V(nir, nir_lower_system_values);
339 NIR_PASS_V(nir, nir_lower_clip_cull_distance_arrays);
340 NIR_PASS_V(nir, radv_nir_lower_ycbcr_textures, layout);
341 }
342
343 /* Vulkan uses the separate-shader linking model */
344 nir->info.separate_shader = true;
345
346 nir_shader_gather_info(nir, nir_shader_get_entrypoint(nir));
347
348 static const nir_lower_tex_options tex_options = {
349 .lower_txp = ~0,
350 .lower_tg4_offsets = true,
351 };
352
353 nir_lower_tex(nir, &tex_options);
354
355 nir_lower_vars_to_ssa(nir);
356
357 if (nir->info.stage == MESA_SHADER_VERTEX ||
358 nir->info.stage == MESA_SHADER_GEOMETRY) {
359 NIR_PASS_V(nir, nir_lower_io_to_temporaries,
360 nir_shader_get_entrypoint(nir), true, true);
361 } else if (nir->info.stage == MESA_SHADER_TESS_EVAL||
362 nir->info.stage == MESA_SHADER_FRAGMENT) {
363 NIR_PASS_V(nir, nir_lower_io_to_temporaries,
364 nir_shader_get_entrypoint(nir), true, false);
365 }
366
367 nir_split_var_copies(nir);
368
369 nir_lower_global_vars_to_local(nir);
370 nir_remove_dead_variables(nir, nir_var_function_temp);
371 nir_lower_subgroups(nir, &(struct nir_lower_subgroups_options) {
372 .subgroup_size = 64,
373 .ballot_bit_size = 64,
374 .lower_to_scalar = 1,
375 .lower_subgroup_masks = 1,
376 .lower_shuffle = 1,
377 .lower_shuffle_to_32bit = 1,
378 .lower_vote_eq_to_ballot = 1,
379 });
380
381 nir_lower_load_const_to_scalar(nir);
382
383 if (!(flags & VK_PIPELINE_CREATE_DISABLE_OPTIMIZATION_BIT))
384 radv_optimize_nir(nir, false, true);
385
386 /* We call nir_lower_var_copies() after the first radv_optimize_nir()
387 * to remove any copies introduced by nir_opt_find_array_copies().
388 */
389 nir_lower_var_copies(nir);
390
391 /* Indirect lowering must be called after the radv_optimize_nir() loop
392 * has been called at least once. Otherwise indirect lowering can
393 * bloat the instruction count of the loop and cause it to be
394 * considered too large for unrolling.
395 */
396 ac_lower_indirect_derefs(nir, device->physical_device->rad_info.chip_class);
397 radv_optimize_nir(nir, flags & VK_PIPELINE_CREATE_DISABLE_OPTIMIZATION_BIT, false);
398
399 return nir;
400 }
401
402 void *
403 radv_alloc_shader_memory(struct radv_device *device,
404 struct radv_shader_variant *shader)
405 {
406 mtx_lock(&device->shader_slab_mutex);
407 list_for_each_entry(struct radv_shader_slab, slab, &device->shader_slabs, slabs) {
408 uint64_t offset = 0;
409 list_for_each_entry(struct radv_shader_variant, s, &slab->shaders, slab_list) {
410 if (s->bo_offset - offset >= shader->code_size) {
411 shader->bo = slab->bo;
412 shader->bo_offset = offset;
413 list_addtail(&shader->slab_list, &s->slab_list);
414 mtx_unlock(&device->shader_slab_mutex);
415 return slab->ptr + offset;
416 }
417 offset = align_u64(s->bo_offset + s->code_size, 256);
418 }
419 if (slab->size - offset >= shader->code_size) {
420 shader->bo = slab->bo;
421 shader->bo_offset = offset;
422 list_addtail(&shader->slab_list, &slab->shaders);
423 mtx_unlock(&device->shader_slab_mutex);
424 return slab->ptr + offset;
425 }
426 }
427
428 mtx_unlock(&device->shader_slab_mutex);
429 struct radv_shader_slab *slab = calloc(1, sizeof(struct radv_shader_slab));
430
431 slab->size = 256 * 1024;
432 slab->bo = device->ws->buffer_create(device->ws, slab->size, 256,
433 RADEON_DOMAIN_VRAM,
434 RADEON_FLAG_NO_INTERPROCESS_SHARING |
435 (device->physical_device->cpdma_prefetch_writes_memory ?
436 0 : RADEON_FLAG_READ_ONLY),
437 RADV_BO_PRIORITY_SHADER);
438 slab->ptr = (char*)device->ws->buffer_map(slab->bo);
439 list_inithead(&slab->shaders);
440
441 mtx_lock(&device->shader_slab_mutex);
442 list_add(&slab->slabs, &device->shader_slabs);
443
444 shader->bo = slab->bo;
445 shader->bo_offset = 0;
446 list_add(&shader->slab_list, &slab->shaders);
447 mtx_unlock(&device->shader_slab_mutex);
448 return slab->ptr;
449 }
450
451 void
452 radv_destroy_shader_slabs(struct radv_device *device)
453 {
454 list_for_each_entry_safe(struct radv_shader_slab, slab, &device->shader_slabs, slabs) {
455 device->ws->buffer_destroy(slab->bo);
456 free(slab);
457 }
458 mtx_destroy(&device->shader_slab_mutex);
459 }
460
461 /* For the UMR disassembler. */
462 #define DEBUGGER_END_OF_CODE_MARKER 0xbf9f0000 /* invalid instruction */
463 #define DEBUGGER_NUM_MARKERS 5
464
465 static unsigned
466 radv_get_shader_binary_size(struct ac_shader_binary *binary)
467 {
468 return binary->code_size + DEBUGGER_NUM_MARKERS * 4;
469 }
470
471 static void
472 radv_fill_shader_variant(struct radv_device *device,
473 struct radv_shader_variant *variant,
474 struct radv_nir_compiler_options *options,
475 struct ac_shader_binary *binary,
476 gl_shader_stage stage)
477 {
478 bool scratch_enabled = variant->config.scratch_bytes_per_wave > 0;
479 struct radv_shader_info *info = &variant->info.info;
480 unsigned vgpr_comp_cnt = 0;
481
482 variant->code_size = radv_get_shader_binary_size(binary);
483 variant->rsrc2 = S_00B12C_USER_SGPR(variant->info.num_user_sgprs) |
484 S_00B12C_USER_SGPR_MSB(variant->info.num_user_sgprs >> 5) |
485 S_00B12C_SCRATCH_EN(scratch_enabled) |
486 S_00B12C_SO_BASE0_EN(!!info->so.strides[0]) |
487 S_00B12C_SO_BASE1_EN(!!info->so.strides[1]) |
488 S_00B12C_SO_BASE2_EN(!!info->so.strides[2]) |
489 S_00B12C_SO_BASE3_EN(!!info->so.strides[3]) |
490 S_00B12C_SO_EN(!!info->so.num_outputs);
491
492 variant->rsrc1 = S_00B848_VGPRS((variant->config.num_vgprs - 1) / 4) |
493 S_00B848_SGPRS((variant->config.num_sgprs - 1) / 8) |
494 S_00B848_DX10_CLAMP(1) |
495 S_00B848_FLOAT_MODE(variant->config.float_mode);
496
497 switch (stage) {
498 case MESA_SHADER_TESS_EVAL:
499 if (options->key.tes.as_es) {
500 assert(device->physical_device->rad_info.chip_class <= GFX8);
501 vgpr_comp_cnt = info->uses_prim_id ? 3 : 2;
502 } else {
503 bool enable_prim_id = options->key.tes.export_prim_id || info->uses_prim_id;
504 vgpr_comp_cnt = enable_prim_id ? 3 : 2;
505 }
506 variant->rsrc2 |= S_00B12C_OC_LDS_EN(1);
507 break;
508 case MESA_SHADER_TESS_CTRL:
509 if (device->physical_device->rad_info.chip_class >= GFX9) {
510 /* We need at least 2 components for LS.
511 * VGPR0-3: (VertexID, RelAutoindex, InstanceID / StepRate0, InstanceID).
512 * StepRate0 is set to 1. so that VGPR3 doesn't have to be loaded.
513 */
514 vgpr_comp_cnt = info->vs.needs_instance_id ? 2 : 1;
515 } else {
516 variant->rsrc2 |= S_00B12C_OC_LDS_EN(1);
517 }
518 break;
519 case MESA_SHADER_VERTEX:
520 if (variant->info.vs.as_ls) {
521 assert(device->physical_device->rad_info.chip_class <= GFX8);
522 /* We need at least 2 components for LS.
523 * VGPR0-3: (VertexID, RelAutoindex, InstanceID / StepRate0, InstanceID).
524 * StepRate0 is set to 1. so that VGPR3 doesn't have to be loaded.
525 */
526 vgpr_comp_cnt = info->vs.needs_instance_id ? 2 : 1;
527 } else if (variant->info.vs.as_es) {
528 assert(device->physical_device->rad_info.chip_class <= GFX8);
529 /* VGPR0-3: (VertexID, InstanceID / StepRate0, ...) */
530 vgpr_comp_cnt = info->vs.needs_instance_id ? 1 : 0;
531 } else {
532 /* VGPR0-3: (VertexID, InstanceID / StepRate0, PrimID, InstanceID)
533 * If PrimID is disabled. InstanceID / StepRate1 is loaded instead.
534 * StepRate0 is set to 1. so that VGPR3 doesn't have to be loaded.
535 */
536 if (options->key.vs.export_prim_id) {
537 vgpr_comp_cnt = 2;
538 } else if (info->vs.needs_instance_id) {
539 vgpr_comp_cnt = 1;
540 } else {
541 vgpr_comp_cnt = 0;
542 }
543 }
544 break;
545 case MESA_SHADER_FRAGMENT:
546 case MESA_SHADER_GEOMETRY:
547 break;
548 case MESA_SHADER_COMPUTE:
549 variant->rsrc2 |=
550 S_00B84C_TGID_X_EN(info->cs.uses_block_id[0]) |
551 S_00B84C_TGID_Y_EN(info->cs.uses_block_id[1]) |
552 S_00B84C_TGID_Z_EN(info->cs.uses_block_id[2]) |
553 S_00B84C_TIDIG_COMP_CNT(info->cs.uses_thread_id[2] ? 2 :
554 info->cs.uses_thread_id[1] ? 1 : 0) |
555 S_00B84C_TG_SIZE_EN(info->cs.uses_local_invocation_idx) |
556 S_00B84C_LDS_SIZE(variant->config.lds_size);
557 break;
558 default:
559 unreachable("unsupported shader type");
560 break;
561 }
562
563 if (device->physical_device->rad_info.chip_class >= GFX9 &&
564 stage == MESA_SHADER_GEOMETRY) {
565 unsigned es_type = variant->info.gs.es_type;
566 unsigned gs_vgpr_comp_cnt, es_vgpr_comp_cnt;
567
568 if (es_type == MESA_SHADER_VERTEX) {
569 /* VGPR0-3: (VertexID, InstanceID / StepRate0, ...) */
570 es_vgpr_comp_cnt = info->vs.needs_instance_id ? 1 : 0;
571 } else if (es_type == MESA_SHADER_TESS_EVAL) {
572 es_vgpr_comp_cnt = info->uses_prim_id ? 3 : 2;
573 } else {
574 unreachable("invalid shader ES type");
575 }
576
577 /* If offsets 4, 5 are used, GS_VGPR_COMP_CNT is ignored and
578 * VGPR[0:4] are always loaded.
579 */
580 if (info->uses_invocation_id) {
581 gs_vgpr_comp_cnt = 3; /* VGPR3 contains InvocationID. */
582 } else if (info->uses_prim_id) {
583 gs_vgpr_comp_cnt = 2; /* VGPR2 contains PrimitiveID. */
584 } else if (variant->info.gs.vertices_in >= 3) {
585 gs_vgpr_comp_cnt = 1; /* VGPR1 contains offsets 2, 3 */
586 } else {
587 gs_vgpr_comp_cnt = 0; /* VGPR0 contains offsets 0, 1 */
588 }
589
590 variant->rsrc1 |= S_00B228_GS_VGPR_COMP_CNT(gs_vgpr_comp_cnt);
591 variant->rsrc2 |= S_00B22C_ES_VGPR_COMP_CNT(es_vgpr_comp_cnt) |
592 S_00B22C_OC_LDS_EN(es_type == MESA_SHADER_TESS_EVAL);
593 } else if (device->physical_device->rad_info.chip_class >= GFX9 &&
594 stage == MESA_SHADER_TESS_CTRL) {
595 variant->rsrc1 |= S_00B428_LS_VGPR_COMP_CNT(vgpr_comp_cnt);
596 } else {
597 variant->rsrc1 |= S_00B128_VGPR_COMP_CNT(vgpr_comp_cnt);
598 }
599
600 void *ptr = radv_alloc_shader_memory(device, variant);
601 memcpy(ptr, binary->code, binary->code_size);
602
603 /* Add end-of-code markers for the UMR disassembler. */
604 uint32_t *ptr32 = (uint32_t *)ptr + binary->code_size / 4;
605 for (unsigned i = 0; i < DEBUGGER_NUM_MARKERS; i++)
606 ptr32[i] = DEBUGGER_END_OF_CODE_MARKER;
607
608 }
609
610 static void radv_init_llvm_target()
611 {
612 LLVMInitializeAMDGPUTargetInfo();
613 LLVMInitializeAMDGPUTarget();
614 LLVMInitializeAMDGPUTargetMC();
615 LLVMInitializeAMDGPUAsmPrinter();
616
617 /* For inline assembly. */
618 LLVMInitializeAMDGPUAsmParser();
619
620 /* Workaround for bug in llvm 4.0 that causes image intrinsics
621 * to disappear.
622 * https://reviews.llvm.org/D26348
623 *
624 * Workaround for bug in llvm that causes the GPU to hang in presence
625 * of nested loops because there is an exec mask issue. The proper
626 * solution is to fix LLVM but this might require a bunch of work.
627 * https://bugs.llvm.org/show_bug.cgi?id=37744
628 *
629 * "mesa" is the prefix for error messages.
630 */
631 if (HAVE_LLVM >= 0x0800) {
632 const char *argv[2] = { "mesa", "-simplifycfg-sink-common=false" };
633 LLVMParseCommandLineOptions(2, argv, NULL);
634
635 } else {
636 const char *argv[3] = { "mesa", "-simplifycfg-sink-common=false",
637 "-amdgpu-skip-threshold=1" };
638 LLVMParseCommandLineOptions(3, argv, NULL);
639 }
640 }
641
642 static once_flag radv_init_llvm_target_once_flag = ONCE_FLAG_INIT;
643
644 static void radv_init_llvm_once(void)
645 {
646 call_once(&radv_init_llvm_target_once_flag, radv_init_llvm_target);
647 }
648
649 static struct radv_shader_variant *
650 shader_variant_create(struct radv_device *device,
651 struct radv_shader_module *module,
652 struct nir_shader * const *shaders,
653 int shader_count,
654 gl_shader_stage stage,
655 struct radv_nir_compiler_options *options,
656 bool gs_copy_shader,
657 void **code_out,
658 unsigned *code_size_out)
659 {
660 enum radeon_family chip_family = device->physical_device->rad_info.family;
661 enum ac_target_machine_options tm_options = 0;
662 struct radv_shader_variant *variant;
663 struct ac_shader_binary binary;
664 struct ac_llvm_compiler ac_llvm;
665 bool thread_compiler;
666 variant = calloc(1, sizeof(struct radv_shader_variant));
667 if (!variant)
668 return NULL;
669
670 options->family = chip_family;
671 options->chip_class = device->physical_device->rad_info.chip_class;
672 options->dump_shader = radv_can_dump_shader(device, module, gs_copy_shader);
673 options->dump_preoptir = options->dump_shader &&
674 device->instance->debug_flags & RADV_DEBUG_PREOPTIR;
675 options->record_llvm_ir = device->keep_shader_info;
676 options->check_ir = device->instance->debug_flags & RADV_DEBUG_CHECKIR;
677 options->tess_offchip_block_dw_size = device->tess_offchip_block_dw_size;
678 options->address32_hi = device->physical_device->rad_info.address32_hi;
679
680 if (options->supports_spill)
681 tm_options |= AC_TM_SUPPORTS_SPILL;
682 if (device->instance->perftest_flags & RADV_PERFTEST_SISCHED)
683 tm_options |= AC_TM_SISCHED;
684 if (options->check_ir)
685 tm_options |= AC_TM_CHECK_IR;
686 if (device->instance->debug_flags & RADV_DEBUG_NO_LOAD_STORE_OPT)
687 tm_options |= AC_TM_NO_LOAD_STORE_OPT;
688
689 thread_compiler = !(device->instance->debug_flags & RADV_DEBUG_NOTHREADLLVM);
690 radv_init_llvm_once();
691 radv_init_llvm_compiler(&ac_llvm,
692 thread_compiler,
693 chip_family, tm_options);
694 if (gs_copy_shader) {
695 assert(shader_count == 1);
696 radv_compile_gs_copy_shader(&ac_llvm, *shaders, &binary,
697 &variant->config, &variant->info,
698 options);
699 } else {
700 radv_compile_nir_shader(&ac_llvm, &binary, &variant->config,
701 &variant->info, shaders, shader_count,
702 options);
703 }
704
705 radv_destroy_llvm_compiler(&ac_llvm, thread_compiler);
706
707 radv_fill_shader_variant(device, variant, options, &binary, stage);
708
709 if (code_out) {
710 *code_out = binary.code;
711 *code_size_out = binary.code_size;
712 } else
713 free(binary.code);
714 free(binary.config);
715 free(binary.rodata);
716 free(binary.global_symbol_offsets);
717 free(binary.relocs);
718 variant->ref_count = 1;
719
720 if (device->keep_shader_info) {
721 variant->disasm_string = binary.disasm_string;
722 variant->llvm_ir_string = binary.llvm_ir_string;
723 if (!gs_copy_shader && !module->nir) {
724 variant->nir = *shaders;
725 variant->spirv = (uint32_t *)module->data;
726 variant->spirv_size = module->size;
727 }
728 } else {
729 free(binary.disasm_string);
730 }
731
732 return variant;
733 }
734
735 struct radv_shader_variant *
736 radv_shader_variant_create(struct radv_device *device,
737 struct radv_shader_module *module,
738 struct nir_shader *const *shaders,
739 int shader_count,
740 struct radv_pipeline_layout *layout,
741 const struct radv_shader_variant_key *key,
742 void **code_out,
743 unsigned *code_size_out)
744 {
745 struct radv_nir_compiler_options options = {0};
746
747 options.layout = layout;
748 if (key)
749 options.key = *key;
750
751 options.unsafe_math = !!(device->instance->debug_flags & RADV_DEBUG_UNSAFE_MATH);
752 options.supports_spill = true;
753
754 return shader_variant_create(device, module, shaders, shader_count, shaders[shader_count - 1]->info.stage,
755 &options, false, code_out, code_size_out);
756 }
757
758 struct radv_shader_variant *
759 radv_create_gs_copy_shader(struct radv_device *device,
760 struct nir_shader *shader,
761 void **code_out,
762 unsigned *code_size_out,
763 bool multiview)
764 {
765 struct radv_nir_compiler_options options = {0};
766
767 options.key.has_multiview_view_index = multiview;
768
769 return shader_variant_create(device, NULL, &shader, 1, MESA_SHADER_VERTEX,
770 &options, true, code_out, code_size_out);
771 }
772
773 void
774 radv_shader_variant_destroy(struct radv_device *device,
775 struct radv_shader_variant *variant)
776 {
777 if (!p_atomic_dec_zero(&variant->ref_count))
778 return;
779
780 mtx_lock(&device->shader_slab_mutex);
781 list_del(&variant->slab_list);
782 mtx_unlock(&device->shader_slab_mutex);
783
784 ralloc_free(variant->nir);
785 free(variant->disasm_string);
786 free(variant->llvm_ir_string);
787 free(variant);
788 }
789
790 const char *
791 radv_get_shader_name(struct radv_shader_variant *var, gl_shader_stage stage)
792 {
793 switch (stage) {
794 case MESA_SHADER_VERTEX: return var->info.vs.as_ls ? "Vertex Shader as LS" : var->info.vs.as_es ? "Vertex Shader as ES" : "Vertex Shader as VS";
795 case MESA_SHADER_GEOMETRY: return "Geometry Shader";
796 case MESA_SHADER_FRAGMENT: return "Pixel Shader";
797 case MESA_SHADER_COMPUTE: return "Compute Shader";
798 case MESA_SHADER_TESS_CTRL: return "Tessellation Control Shader";
799 case MESA_SHADER_TESS_EVAL: return var->info.tes.as_es ? "Tessellation Evaluation Shader as ES" : "Tessellation Evaluation Shader as VS";
800 default:
801 return "Unknown shader";
802 };
803 }
804
805 static void
806 generate_shader_stats(struct radv_device *device,
807 struct radv_shader_variant *variant,
808 gl_shader_stage stage,
809 struct _mesa_string_buffer *buf)
810 {
811 enum chip_class chip_class = device->physical_device->rad_info.chip_class;
812 unsigned lds_increment = chip_class >= GFX7 ? 512 : 256;
813 struct ac_shader_config *conf;
814 unsigned max_simd_waves;
815 unsigned lds_per_wave = 0;
816
817 max_simd_waves = ac_get_max_simd_waves(device->physical_device->rad_info.family);
818
819 conf = &variant->config;
820
821 if (stage == MESA_SHADER_FRAGMENT) {
822 lds_per_wave = conf->lds_size * lds_increment +
823 align(variant->info.fs.num_interp * 48,
824 lds_increment);
825 } else if (stage == MESA_SHADER_COMPUTE) {
826 unsigned max_workgroup_size =
827 radv_nir_get_max_workgroup_size(chip_class, variant->nir);
828 lds_per_wave = (conf->lds_size * lds_increment) /
829 DIV_ROUND_UP(max_workgroup_size, 64);
830 }
831
832 if (conf->num_sgprs)
833 max_simd_waves =
834 MIN2(max_simd_waves,
835 ac_get_num_physical_sgprs(chip_class) / conf->num_sgprs);
836
837 if (conf->num_vgprs)
838 max_simd_waves =
839 MIN2(max_simd_waves,
840 RADV_NUM_PHYSICAL_VGPRS / conf->num_vgprs);
841
842 /* LDS is 64KB per CU (4 SIMDs), divided into 16KB blocks per SIMD
843 * that PS can use.
844 */
845 if (lds_per_wave)
846 max_simd_waves = MIN2(max_simd_waves, 16384 / lds_per_wave);
847
848 if (stage == MESA_SHADER_FRAGMENT) {
849 _mesa_string_buffer_printf(buf, "*** SHADER CONFIG ***\n"
850 "SPI_PS_INPUT_ADDR = 0x%04x\n"
851 "SPI_PS_INPUT_ENA = 0x%04x\n",
852 conf->spi_ps_input_addr, conf->spi_ps_input_ena);
853 }
854
855 _mesa_string_buffer_printf(buf, "*** SHADER STATS ***\n"
856 "SGPRS: %d\n"
857 "VGPRS: %d\n"
858 "Spilled SGPRs: %d\n"
859 "Spilled VGPRs: %d\n"
860 "PrivMem VGPRS: %d\n"
861 "Code Size: %d bytes\n"
862 "LDS: %d blocks\n"
863 "Scratch: %d bytes per wave\n"
864 "Max Waves: %d\n"
865 "********************\n\n\n",
866 conf->num_sgprs, conf->num_vgprs,
867 conf->spilled_sgprs, conf->spilled_vgprs,
868 variant->info.private_mem_vgprs, variant->code_size,
869 conf->lds_size, conf->scratch_bytes_per_wave,
870 max_simd_waves);
871 }
872
873 void
874 radv_shader_dump_stats(struct radv_device *device,
875 struct radv_shader_variant *variant,
876 gl_shader_stage stage,
877 FILE *file)
878 {
879 struct _mesa_string_buffer *buf = _mesa_string_buffer_create(NULL, 256);
880
881 generate_shader_stats(device, variant, stage, buf);
882
883 fprintf(file, "\n%s:\n", radv_get_shader_name(variant, stage));
884 fprintf(file, "%s", buf->buf);
885
886 _mesa_string_buffer_destroy(buf);
887 }
888
889 VkResult
890 radv_GetShaderInfoAMD(VkDevice _device,
891 VkPipeline _pipeline,
892 VkShaderStageFlagBits shaderStage,
893 VkShaderInfoTypeAMD infoType,
894 size_t* pInfoSize,
895 void* pInfo)
896 {
897 RADV_FROM_HANDLE(radv_device, device, _device);
898 RADV_FROM_HANDLE(radv_pipeline, pipeline, _pipeline);
899 gl_shader_stage stage = vk_to_mesa_shader_stage(shaderStage);
900 struct radv_shader_variant *variant = pipeline->shaders[stage];
901 struct _mesa_string_buffer *buf;
902 VkResult result = VK_SUCCESS;
903
904 /* Spec doesn't indicate what to do if the stage is invalid, so just
905 * return no info for this. */
906 if (!variant)
907 return vk_error(device->instance, VK_ERROR_FEATURE_NOT_PRESENT);
908
909 switch (infoType) {
910 case VK_SHADER_INFO_TYPE_STATISTICS_AMD:
911 if (!pInfo) {
912 *pInfoSize = sizeof(VkShaderStatisticsInfoAMD);
913 } else {
914 unsigned lds_multiplier = device->physical_device->rad_info.chip_class >= GFX7 ? 512 : 256;
915 struct ac_shader_config *conf = &variant->config;
916
917 VkShaderStatisticsInfoAMD statistics = {};
918 statistics.shaderStageMask = shaderStage;
919 statistics.numPhysicalVgprs = RADV_NUM_PHYSICAL_VGPRS;
920 statistics.numPhysicalSgprs = ac_get_num_physical_sgprs(device->physical_device->rad_info.chip_class);
921 statistics.numAvailableSgprs = statistics.numPhysicalSgprs;
922
923 if (stage == MESA_SHADER_COMPUTE) {
924 unsigned *local_size = variant->nir->info.cs.local_size;
925 unsigned workgroup_size = local_size[0] * local_size[1] * local_size[2];
926
927 statistics.numAvailableVgprs = statistics.numPhysicalVgprs /
928 ceil((double)workgroup_size / statistics.numPhysicalVgprs);
929
930 statistics.computeWorkGroupSize[0] = local_size[0];
931 statistics.computeWorkGroupSize[1] = local_size[1];
932 statistics.computeWorkGroupSize[2] = local_size[2];
933 } else {
934 statistics.numAvailableVgprs = statistics.numPhysicalVgprs;
935 }
936
937 statistics.resourceUsage.numUsedVgprs = conf->num_vgprs;
938 statistics.resourceUsage.numUsedSgprs = conf->num_sgprs;
939 statistics.resourceUsage.ldsSizePerLocalWorkGroup = 32768;
940 statistics.resourceUsage.ldsUsageSizeInBytes = conf->lds_size * lds_multiplier;
941 statistics.resourceUsage.scratchMemUsageInBytes = conf->scratch_bytes_per_wave;
942
943 size_t size = *pInfoSize;
944 *pInfoSize = sizeof(statistics);
945
946 memcpy(pInfo, &statistics, MIN2(size, *pInfoSize));
947
948 if (size < *pInfoSize)
949 result = VK_INCOMPLETE;
950 }
951
952 break;
953 case VK_SHADER_INFO_TYPE_DISASSEMBLY_AMD:
954 buf = _mesa_string_buffer_create(NULL, 1024);
955
956 _mesa_string_buffer_printf(buf, "%s:\n", radv_get_shader_name(variant, stage));
957 _mesa_string_buffer_printf(buf, "%s\n\n", variant->llvm_ir_string);
958 _mesa_string_buffer_printf(buf, "%s\n\n", variant->disasm_string);
959 generate_shader_stats(device, variant, stage, buf);
960
961 /* Need to include the null terminator. */
962 size_t length = buf->length + 1;
963
964 if (!pInfo) {
965 *pInfoSize = length;
966 } else {
967 size_t size = *pInfoSize;
968 *pInfoSize = length;
969
970 memcpy(pInfo, buf->buf, MIN2(size, length));
971
972 if (size < length)
973 result = VK_INCOMPLETE;
974 }
975
976 _mesa_string_buffer_destroy(buf);
977 break;
978 default:
979 /* VK_SHADER_INFO_TYPE_BINARY_AMD unimplemented for now. */
980 result = VK_ERROR_FEATURE_NOT_PRESENT;
981 break;
982 }
983
984 return result;
985 }