diff --git a/src/amd/vulkan/radv_pipeline_rt.c b/src/amd/vulkan/radv_pipeline_rt.c index 745ef8a59af..dc6b6c47466 100644 --- a/src/amd/vulkan/radv_pipeline_rt.c +++ b/src/amd/vulkan/radv_pipeline_rt.c @@ -22,7 +22,6 @@ */ #include "nir/nir.h" -#include "nir/nir_control_flow.h" #include "radv_debug.h" #include "radv_private.h" diff --git a/src/amd/vulkan/radv_rt_shader.c b/src/amd/vulkan/radv_rt_shader.c index 8bc0e100f95..4eb4a2d21b1 100644 --- a/src/amd/vulkan/radv_rt_shader.c +++ b/src/amd/vulkan/radv_rt_shader.c @@ -1588,132 +1588,6 @@ radv_build_traversal_shader(struct radv_device *device, struct radv_ray_tracing_ return b.shader; } -static bool -should_move_rt_instruction(nir_intrinsic_op intrinsic) -{ - switch (intrinsic) { - case nir_intrinsic_load_hit_attrib_amd: - case nir_intrinsic_load_rt_arg_scratch_offset_amd: - case nir_intrinsic_load_ray_flags: - case nir_intrinsic_load_ray_object_origin: - case nir_intrinsic_load_ray_world_origin: - case nir_intrinsic_load_ray_t_min: - case nir_intrinsic_load_ray_object_direction: - case nir_intrinsic_load_ray_world_direction: - case nir_intrinsic_load_ray_t_max: - return true; - default: - return false; - } -} - -static void -move_rt_instructions(nir_shader *shader) -{ - nir_cursor target = nir_before_cf_list(&nir_shader_get_entrypoint(shader)->body); - - nir_foreach_block (block, nir_shader_get_entrypoint(shader)) { - nir_foreach_instr_safe (instr, block) { - if (instr->type != nir_instr_type_intrinsic) - continue; - - nir_intrinsic_instr *intrinsic = nir_instr_as_intrinsic(instr); - - if (!should_move_rt_instruction(intrinsic->intrinsic)) - continue; - - nir_instr_move(target, instr); - } - } - - nir_metadata_preserve(nir_shader_get_entrypoint(shader), - nir_metadata_all & (~nir_metadata_instr_index)); -} - -nir_shader * -create_rt_shader(struct radv_device *device, struct radv_ray_tracing_pipeline *pipeline, - const VkRayTracingPipelineCreateInfoKHR *pCreateInfo, - const struct radv_pipeline_key *key) -{ - nir_builder b = radv_meta_init_shader(device, MESA_SHADER_RAYGEN, "rt_combined"); - b.shader->info.internal = false; - b.shader->info.workgroup_size[0] = 8; - b.shader->info.workgroup_size[1] = device->physical_device->rt_wave_size == 64 ? 8 : 4; - b.shader->info.shared_size = device->physical_device->rt_wave_size * RADV_MAX_HIT_ATTRIB_SIZE; - - struct rt_variables vars = create_rt_variables(b.shader, pCreateInfo->flags); - load_sbt_entry(&b, &vars, nir_imm_int(&b, 0), SBT_RAYGEN, SBT_GENERAL_IDX); - nir_store_var(&b, vars.stack_ptr, nir_load_rt_dynamic_callable_stack_base_amd(&b), 0x1); - - nir_loop *loop = nir_push_loop(&b); - nir_ssa_def *idx = nir_load_var(&b, vars.idx); - - /* Insert traversal shader */ - nir_shader *traversal = radv_build_traversal_shader(device, pipeline, pCreateInfo, key); - b.shader->info.shared_size = MAX2(b.shader->info.shared_size, traversal->info.shared_size); - assert(b.shader->info.shared_size <= 32768); - insert_rt_case(&b, traversal, &vars, idx, 0, 1, -1u, NULL); - ralloc_free(traversal); - - struct radv_ray_tracing_group *groups = pipeline->groups; - struct radv_ray_tracing_stage *stages = pipeline->stages; - unsigned call_idx_base = 1; - for (unsigned i = 0; i < pCreateInfo->groupCount; ++i) { - unsigned stage_idx = groups[i].recursive_shader; - - if (stage_idx == VK_SHADER_UNUSED_KHR) - continue; - - /* Avoid emitting stages with the same shaders/handles multiple times. */ - bool is_dup = false; - for (unsigned j = 0; j < i; ++j) - if (groups[j].handle.general_index == groups[i].handle.general_index) - is_dup = true; - - if (is_dup) - continue; - - nir_shader *nir_stage = radv_pipeline_cache_handle_to_nir(device, stages[stage_idx].shader); - assert(nir_stage); - ASSERTED gl_shader_stage type = nir_stage->info.stage; - assert(type == MESA_SHADER_RAYGEN || type == MESA_SHADER_CALLABLE || - type == MESA_SHADER_CLOSEST_HIT || type == MESA_SHADER_MISS); - - /* Move ray tracing system values to the top that are set by rt_trace_ray - * to prevent them from being overwritten by other rt_trace_ray calls. - */ - NIR_PASS_V(nir_stage, move_rt_instructions); - - const nir_lower_shader_calls_options opts = { - .address_format = nir_address_format_32bit_offset, - .stack_alignment = 16, - .localized_loads = true, - .vectorizer_callback = radv_mem_vectorize_callback, - }; - uint32_t num_resume_shaders = 0; - nir_shader **resume_shaders = NULL; - nir_lower_shader_calls(nir_stage, &opts, &resume_shaders, &num_resume_shaders, nir_stage); - - insert_rt_case(&b, nir_stage, &vars, idx, call_idx_base, groups[i].handle.general_index, - stage_idx, stages); - for (unsigned j = 0; j < num_resume_shaders; ++j) { - insert_rt_case(&b, resume_shaders[j], &vars, idx, call_idx_base, call_idx_base + 1 + j, - stage_idx, stages); - } - - ralloc_free(nir_stage); - call_idx_base += num_resume_shaders; - } - - nir_pop_loop(&b, loop); - - /* Deal with all the inline functions. */ - nir_index_ssa_defs(nir_shader_get_entrypoint(b.shader)); - nir_metadata_preserve(nir_shader_get_entrypoint(b.shader), nir_metadata_none); - - return b.shader; -} - /** Select the next shader based on priorities: * * Detect the priority of the shader stage by the lowest bits in the address (low to high): diff --git a/src/amd/vulkan/radv_shader.h b/src/amd/vulkan/radv_shader.h index 0fc89fae7e8..c2e99e77b35 100644 --- a/src/amd/vulkan/radv_shader.h +++ b/src/amd/vulkan/radv_shader.h @@ -786,10 +786,6 @@ bool radv_consider_culling(const struct radv_physical_device *pdevice, struct ni void radv_get_nir_options(struct radv_physical_device *device); -nir_shader *create_rt_shader(struct radv_device *device, struct radv_ray_tracing_pipeline *pipeline, - const VkRayTracingPipelineCreateInfoKHR *pCreateInfo, - const struct radv_pipeline_key *key); - nir_shader *radv_build_traversal_shader(struct radv_device *device, struct radv_ray_tracing_pipeline *pipeline, const VkRayTracingPipelineCreateInfoKHR *pCreateInfo,