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):