diff --git a/src/amd/vulkan/radv_pipeline_compute.c b/src/amd/vulkan/radv_pipeline_compute.c index 477304baabf..f6836f0cd22 100644 --- a/src/amd/vulkan/radv_pipeline_compute.c +++ b/src/amd/vulkan/radv_pipeline_compute.c @@ -105,7 +105,7 @@ radv_compute_pipeline_init(struct radv_compute_pipeline *pipeline, const struct struct radv_shader * radv_compile_cs(struct radv_device *device, struct vk_pipeline_cache *cache, struct radv_shader_stage *cs_stage, - bool keep_executable_info, bool keep_statistic_info, bool is_internal, + bool keep_executable_info, bool keep_statistic_info, bool is_internal, bool is_indirect_bindable, struct radv_shader_binary **cs_binary) { struct radv_shader *cs_shader; @@ -121,7 +121,7 @@ radv_compile_cs(struct radv_device *device, struct vk_pipeline_cache *cache, str /* Run the shader info pass. */ radv_nir_shader_info_init(cs_stage->stage, MESA_SHADER_NONE, &cs_stage->info); radv_nir_shader_info_pass(device, cs_stage->nir, &cs_stage->layout, &cs_stage->key, NULL, RADV_PIPELINE_COMPUTE, - false, &cs_stage->info); + false, is_indirect_bindable, &cs_stage->info); radv_declare_shader_args(device, NULL, &cs_stage->info, MESA_SHADER_COMPUTE, MESA_SHADER_NONE, &cs_stage->args); @@ -217,11 +217,13 @@ radv_compute_pipeline_compile(const VkComputePipelineCreateInfo *pCreateInfo, st const struct radv_shader_stage_key stage_key = radv_pipeline_get_shader_key(device, &pCreateInfo->stage, pipeline->base.create_flags, pCreateInfo->pNext); + const bool is_indirect_bindable = !!(pipeline->base.create_flags & VK_PIPELINE_CREATE_INDIRECT_BINDABLE_BIT_NV); radv_pipeline_stage_init(pStage, pipeline_layout, &stage_key, &cs_stage); - pipeline->base.shaders[MESA_SHADER_COMPUTE] = radv_compile_cs( - device, cache, &cs_stage, keep_executable_info, keep_statistic_info, pipeline->base.is_internal, &cs_binary); + pipeline->base.shaders[MESA_SHADER_COMPUTE] = + radv_compile_cs(device, cache, &cs_stage, keep_executable_info, keep_statistic_info, pipeline->base.is_internal, + is_indirect_bindable, &cs_binary); cs_stage.feedback.duration += os_time_get_nano() - stage_start; diff --git a/src/amd/vulkan/radv_pipeline_compute.h b/src/amd/vulkan/radv_pipeline_compute.h index 8d26a8756af..4a43f30ec23 100644 --- a/src/amd/vulkan/radv_pipeline_compute.h +++ b/src/amd/vulkan/radv_pipeline_compute.h @@ -53,7 +53,8 @@ void radv_compute_pipeline_init(struct radv_compute_pipeline *pipeline, const st struct radv_shader *radv_compile_cs(struct radv_device *device, struct vk_pipeline_cache *cache, struct radv_shader_stage *cs_stage, bool keep_executable_info, - bool keep_statistic_info, bool is_internal, struct radv_shader_binary **cs_binary); + bool keep_statistic_info, bool is_internal, bool is_indirect_bindable, + struct radv_shader_binary **cs_binary); VkResult radv_compute_pipeline_create(VkDevice _device, VkPipelineCache _cache, const VkComputePipelineCreateInfo *pCreateInfo, diff --git a/src/amd/vulkan/radv_pipeline_graphics.c b/src/amd/vulkan/radv_pipeline_graphics.c index 3614bc13a47..803a8ac1c50 100644 --- a/src/amd/vulkan/radv_pipeline_graphics.c +++ b/src/amd/vulkan/radv_pipeline_graphics.c @@ -2066,7 +2066,7 @@ radv_fill_shader_info(struct radv_device *device, const enum radv_pipeline_type } radv_nir_shader_info_pass(device, stages[i].nir, &stages[i].layout, &stages[i].key, gfx_state, pipeline_type, - consider_force_vrs, &stages[i].info); + consider_force_vrs, false, &stages[i].info); } radv_nir_shader_info_link(device, gfx_state, stages); @@ -2147,7 +2147,7 @@ radv_create_gs_copy_shader(struct radv_device *device, struct vk_pipeline_cache }; radv_nir_shader_info_init(gs_copy_stage.stage, MESA_SHADER_FRAGMENT, &gs_copy_stage.info); radv_nir_shader_info_pass(device, nir, &gs_stage->layout, &gs_stage->key, gfx_state, RADV_PIPELINE_GRAPHICS, false, - &gs_copy_stage.info); + false, &gs_copy_stage.info); gs_copy_stage.info.wave_size = 64; /* Wave32 not supported. */ gs_copy_stage.info.workgroup_size = 64; /* HW VS: separate waves, no workgroups */ gs_copy_stage.info.so = gs_info->so; diff --git a/src/amd/vulkan/radv_pipeline_rt.c b/src/amd/vulkan/radv_pipeline_rt.c index 2a69006d70d..14b2e14ff4d 100644 --- a/src/amd/vulkan/radv_pipeline_rt.c +++ b/src/amd/vulkan/radv_pipeline_rt.c @@ -361,7 +361,7 @@ radv_rt_nir_to_asm(struct radv_device *device, struct vk_pipeline_cache *cache, nir_shader_gather_info(stage->nir, nir_shader_get_entrypoint(stage->nir)); radv_nir_shader_info_init(stage->stage, MESA_SHADER_NONE, &stage->info); radv_nir_shader_info_pass(device, stage->nir, &stage->layout, &stage->key, NULL, RADV_PIPELINE_RAY_TRACING, false, - &stage->info); + false, &stage->info); /* Declare shader arguments. */ radv_declare_shader_args(device, NULL, &stage->info, stage->stage, MESA_SHADER_NONE, &stage->args); diff --git a/src/amd/vulkan/radv_shader_args.c b/src/amd/vulkan/radv_shader_args.c index 60d06f11bd6..10f1ee3ad07 100644 --- a/src/amd/vulkan/radv_shader_args.c +++ b/src/amd/vulkan/radv_shader_args.c @@ -899,7 +899,7 @@ radv_declare_shader_args(const struct radv_device *device, const struct radv_gra uint32_t num_desc_set = util_bitcount(info->desc_set_used_mask); - if (info->merged_shader_compiled_separately || remaining_sgprs < num_desc_set) { + if (info->force_indirect_desc_sets || remaining_sgprs < num_desc_set) { user_sgpr_info.indirect_all_descriptor_sets = true; user_sgpr_info.remaining_sgprs--; } else { diff --git a/src/amd/vulkan/radv_shader_info.c b/src/amd/vulkan/radv_shader_info.c index 07086173e11..3e85d018dd5 100644 --- a/src/amd/vulkan/radv_shader_info.c +++ b/src/amd/vulkan/radv_shader_info.c @@ -1121,7 +1121,7 @@ void radv_nir_shader_info_pass(struct radv_device *device, const struct nir_shader *nir, const struct radv_shader_layout *layout, const struct radv_shader_stage_key *stage_key, const struct radv_graphics_state_key *gfx_state, const enum radv_pipeline_type pipeline_type, - bool consider_force_vrs, struct radv_shader_info *info) + bool consider_force_vrs, bool is_indirect_bindable, struct radv_shader_info *info) { const struct radv_physical_device *pdev = radv_device_physical(device); struct nir_function *func = (struct nir_function *)exec_list_get_head_const(&nir->functions); @@ -1232,6 +1232,7 @@ radv_nir_shader_info_pass(struct radv_device *device, const struct nir_shader *n info->user_data_0 = radv_get_user_data_0(device, info); info->merged_shader_compiled_separately = radv_is_merged_shader_compiled_separately(device, info); + info->force_indirect_desc_sets = info->merged_shader_compiled_separately || is_indirect_bindable; switch (nir->info.stage) { case MESA_SHADER_COMPUTE: diff --git a/src/amd/vulkan/radv_shader_info.h b/src/amd/vulkan/radv_shader_info.h index 5f91056b1f6..781c0717caf 100644 --- a/src/amd/vulkan/radv_shader_info.h +++ b/src/amd/vulkan/radv_shader_info.h @@ -108,6 +108,7 @@ struct radv_shader_info { bool outputs_linked; bool has_epilog; /* Only for TCS or PS */ bool merged_shader_compiled_separately; /* GFX9+ */ + bool force_indirect_desc_sets; struct { uint8_t output_usage_mask[VARYING_SLOT_VAR31 + 1]; @@ -316,7 +317,7 @@ void radv_nir_shader_info_pass(struct radv_device *device, const struct nir_shad const struct radv_shader_layout *layout, const struct radv_shader_stage_key *stage_key, const struct radv_graphics_state_key *gfx_state, const enum radv_pipeline_type pipeline_type, bool consider_force_vrs, - struct radv_shader_info *info); + bool is_indirect_bindable, struct radv_shader_info *info); void gfx10_get_ngg_info(const struct radv_device *device, struct radv_shader_info *es_info, struct radv_shader_info *gs_info, struct gfx10_ngg_info *out); diff --git a/src/amd/vulkan/radv_shader_object.c b/src/amd/vulkan/radv_shader_object.c index e36196a6dc9..db6004719ce 100644 --- a/src/amd/vulkan/radv_shader_object.c +++ b/src/amd/vulkan/radv_shader_object.c @@ -248,7 +248,7 @@ radv_shader_object_init_compute(struct radv_shader_object *shader_obj, struct ra radv_shader_stage_init(pCreateInfo, &stage); - struct radv_shader *cs_shader = radv_compile_cs(device, NULL, &stage, true, false, false, &cs_binary); + struct radv_shader *cs_shader = radv_compile_cs(device, NULL, &stage, true, false, false, false, &cs_binary); ralloc_free(stage.nir);