radv: fix dumping SPIR-V into hang reports

[mesa.git] / src / amd / vulkan / radv_shader.c

diff --git a/src/amd/vulkan/radv_shader.c b/src/amd/vulkan/radv_shader.c
index 09c5ce639c8734a2781603ebf970d4f060e86ac6..c841a2f072623aab55372437d716eaf39f654d44 100644 (file)
--- a/src/amd/vulkan/radv_shader.c
+++ b/src/amd/vulkan/radv_shader.c
@@ -91,6 +91,7 @@ static const struct nir_shader_compiler_options nir_options_aco = {
        .lower_flrp64 = true,
        .lower_device_index_to_zero = true,
        .lower_fdiv = true,
+       .lower_fmod = true,
        .lower_bitfield_insert_to_bitfield_select = true,
        .lower_bitfield_extract = true,
        .lower_pack_snorm_2x16 = true,
@@ -311,7 +312,7 @@ radv_shader_compile_to_nir(struct radv_device *device,
                assert(module->size % 4 == 0);
 
                if (device->instance->debug_flags & RADV_DEBUG_DUMP_SPIRV)
-                       radv_print_spirv(spirv, module->size, stderr);
+                       radv_print_spirv(module->data, module->size, stderr);
 
                uint32_t num_spec_entries = 0;
                struct nir_spirv_specialization *spec_entries = NULL;
@@ -343,6 +344,7 @@ radv_shader_compile_to_nir(struct radv_device *device,
                                .descriptor_indexing = true,
                                .device_group = true,
                                .draw_parameters = true,
+                               .float_controls = true,
                                .float16 = !device->physical_device->use_aco,
                                .float64 = true,
                                .geometry_streams = true,
@@ -356,6 +358,7 @@ radv_shader_compile_to_nir(struct radv_device *device,
                                .physical_storage_buffer_address = true,
                                .post_depth_coverage = true,
                                .runtime_descriptor_array = true,
+                               .shader_clock = true,
                                .shader_viewport_index_layer = true,
                                .stencil_export = true,
                                .storage_8bit = !device->physical_device->use_aco,
@@ -970,7 +973,14 @@ radv_shader_variant_create(struct radv_device *device,
        variant->info = binary->info;
        radv_postprocess_config(device->physical_device, &config, &binary->info,
                                binary->stage, &variant->config);
-       
+
+       if (radv_device_use_secure_compile(device->instance)) {
+               if (binary->type == RADV_BINARY_TYPE_RTLD)
+                       ac_rtld_close(&rtld_binary);
+
+               return variant;
+       }
+
        void *dest_ptr = radv_alloc_shader_memory(device, variant);
 
        if (binary->type == RADV_BINARY_TYPE_RTLD) {
@@ -1135,7 +1145,14 @@ shader_variant_compile(struct radv_device *device,
        if (keep_shader_info) {
                variant->nir_string = radv_dump_nir_shaders(shaders, shader_count);
                if (!gs_copy_shader && !module->nir) {
-                       variant->spirv = (uint32_t *)module->data;
+                       variant->spirv = malloc(module->size);
+                       if (!variant->spirv) {
+                               free(variant);
+                               free(binary);
+                               return NULL;
+                       }
+
+                       memcpy(variant->spirv, module->data, module->size);
                        variant->spirv_size = module->size;
                }
        }
@@ -1201,6 +1218,7 @@ radv_shader_variant_destroy(struct radv_device *device,
        list_del(&variant->slab_list);
        mtx_unlock(&device->shader_slab_mutex);
 
+       free(variant->spirv);
        free(variant->nir_string);
        free(variant->disasm_string);
        free(variant->ir_string);
@@ -1286,16 +1304,20 @@ radv_get_max_waves(struct radv_device *device,
                               DIV_ROUND_UP(max_workgroup_size, wave_size);
        }
 
-       if (conf->num_sgprs)
+       if (conf->num_sgprs) {
+               unsigned sgprs = align(conf->num_sgprs, chip_class >= GFX8 ? 16 : 8);
                max_simd_waves =
                        MIN2(max_simd_waves,
                             device->physical_device->rad_info.num_physical_sgprs_per_simd /
-                            conf->num_sgprs);
+                            sgprs);
+       }
 
-       if (conf->num_vgprs)
+       if (conf->num_vgprs) {
+               unsigned vgprs = align(conf->num_vgprs, wave_size == 32 ? 8 : 4);
                max_simd_waves =
                        MIN2(max_simd_waves,
-                            RADV_NUM_PHYSICAL_VGPRS / conf->num_vgprs);
+                            RADV_NUM_PHYSICAL_VGPRS / vgprs);
+       }
 
        /* LDS is 64KB per CU (4 SIMDs), divided into 16KB blocks per SIMD
         * that PS can use.