OpenHarmony/third_party_mesa3d
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Revert "radv/rt: Lower ray payloads to registers"

Browse Source 
This causes GPU hangs with vkd3d-proton in CI.

This reverts commit 658ce711d5.

Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/26983>
This commit is contained in:
Samuel Pitoiset
2024-01-10 18:25:34 +01:00
committed by Marge Bot
parent a077c14f15
commit 59c05b9cfa
3 changed files with 26 additions and 72 deletions
Download Patch File Download Diff File Expand all files Collapse all files

87
src/amd/vulkan/nir/radv_nir_rt_shader.c
View File

@@ -187,7 +187,6 @@ lower_rt_derefs(nir_shader *shader)
struct rt_variables {
struct radv_device *device;
const VkPipelineCreateFlags2KHR flags;
bool monolithic;
/* idx of the next shader to run in the next iteration of the main loop.
* During traversal, idx is used to store the SBT index and will contain
@@ -231,13 +230,11 @@ struct rt_variables {
};
static struct rt_variables
create_rt_variables(nir_shader *shader, struct radv_device *device, const VkPipelineCreateFlags2KHR flags,
bool monolithic)
create_rt_variables(nir_shader *shader, struct radv_device *device, const VkPipelineCreateFlags2KHR flags)
{
struct rt_variables vars = {
.device = device,
.flags = flags,
.monolithic = monolithic,
};
vars.idx = nir_variable_create(shader, nir_var_shader_temp, glsl_uint_type(), "idx");
vars.shader_addr = nir_variable_create(shader, nir_var_shader_temp, glsl_uint64_t_type(), "shader_addr");
@@ -793,7 +790,7 @@ insert_rt_case(nir_builder *b, nir_shader *shader, struct rt_variables *vars, ni
nir_opt_dead_cf(shader);
struct rt_variables src_vars = create_rt_variables(shader, vars->device, vars->flags, vars->monolithic);
struct rt_variables src_vars = create_rt_variables(shader, vars->device, vars->flags);
map_rt_variables(var_remap, &src_vars, vars);
NIR_PASS_V(shader, lower_rt_instructions, &src_vars, false);
@@ -803,23 +800,6 @@ insert_rt_case(nir_builder *b, nir_shader *shader, struct rt_variables *vars, ni
inline_constants(b->shader, shader);
if (vars->monolithic) {
/* Work around vulkancts declaring an incompatible dummy payload. */
nir_remove_dead_variables(shader, nir_var_shader_call_data, NULL);
nir_foreach_variable_in_shader (var, shader) {
if (var->data.mode != nir_var_shader_call_data)
continue;
/* There can only be one shader_call_data variable which has to match the caller payload. */
if (var->type != vars->arg->type)
return;
_mesa_hash_table_insert(var_remap, var, vars->arg);
break;
}
}
nir_push_if(b, nir_ieq_imm(b, idx, call_idx));
nir_inline_function_impl(b, nir_shader_get_entrypoint(shader), NULL, var_remap);
nir_pop_if(b, NULL);
@@ -827,17 +807,25 @@ insert_rt_case(nir_builder *b, nir_shader *shader, struct rt_variables *vars, ni
ralloc_free(var_remap);
}
void
radv_nir_lower_rt_io(nir_shader *nir, bool monolithic)
nir_shader *
radv_parse_rt_stage(struct radv_device *device, const VkPipelineShaderStageCreateInfo *sinfo,
const struct radv_pipeline_key *key, const struct radv_pipeline_layout *pipeline_layout)
{
if (!monolithic) {
NIR_PASS(_, nir, nir_lower_vars_to_explicit_types, nir_var_function_temp | nir_var_shader_call_data,
glsl_get_natural_size_align_bytes);
struct radv_shader_stage rt_stage;
NIR_PASS(_, nir, lower_rt_derefs);
radv_pipeline_stage_init(sinfo, pipeline_layout, &rt_stage);
NIR_PASS(_, nir, nir_lower_explicit_io, nir_var_function_temp, nir_address_format_32bit_offset);
}
nir_shader *shader = radv_shader_spirv_to_nir(device, &rt_stage, key, false);
NIR_PASS(_, shader, nir_lower_vars_to_explicit_types, nir_var_function_temp | nir_var_shader_call_data,
glsl_get_natural_size_align_bytes);
NIR_PASS(_, shader, lower_rt_derefs);
NIR_PASS(_, shader, radv_nir_lower_hit_attrib_derefs);
NIR_PASS(_, shader, nir_lower_explicit_io, nir_var_function_temp, nir_address_format_32bit_offset);
return shader;
}
static nir_function_impl *
@@ -1135,8 +1123,6 @@ radv_build_ahit_case(nir_builder *b, nir_def *sbt_idx, struct radv_ray_tracing_g
radv_pipeline_cache_handle_to_nir(data->device, data->pipeline->stages[group->any_hit_shader].nir);
assert(nir_stage);
radv_nir_lower_rt_io(nir_stage, data->vars->monolithic);
insert_rt_case(b, nir_stage, data->vars, sbt_idx, group->handle.any_hit_index);
ralloc_free(nir_stage);
}
@@ -1159,16 +1145,12 @@ radv_build_isec_case(nir_builder *b, nir_def *sbt_idx, struct radv_ray_tracing_g
radv_pipeline_cache_handle_to_nir(data->device, data->pipeline->stages[group->intersection_shader].nir);
assert(nir_stage);
radv_nir_lower_rt_io(nir_stage, data->vars->monolithic);
nir_shader *any_hit_stage = NULL;
if (group->any_hit_shader != VK_SHADER_UNUSED_KHR) {
any_hit_stage =
radv_pipeline_cache_handle_to_nir(data->device, data->pipeline->stages[group->any_hit_shader].nir);
assert(any_hit_stage);
radv_nir_lower_rt_io(any_hit_stage, data->vars->monolithic);
/* reserve stack size for any_hit before it is inlined */
data->pipeline->stages[group->any_hit_shader].stack_size = any_hit_stage->scratch_size;
@@ -1211,8 +1193,6 @@ radv_build_recursive_case(nir_builder *b, nir_def *sbt_idx, struct radv_ray_trac
radv_pipeline_cache_handle_to_nir(data->device, data->pipeline->stages[group->recursive_shader].nir);
assert(nir_stage);
radv_nir_lower_rt_io(nir_stage, data->vars->monolithic);
insert_rt_case(b, nir_stage, data->vars, sbt_idx, group->handle.general_index);
ralloc_free(nir_stage);
}
@@ -1534,7 +1514,7 @@ radv_build_traversal_shader(struct radv_device *device, struct radv_ray_tracing_
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 * MAX_STACK_ENTRY_COUNT * sizeof(uint32_t);
struct rt_variables vars = create_rt_variables(b.shader, device, create_flags, false);
struct rt_variables vars = create_rt_variables(b.shader, device, create_flags);
/* initialize trace_ray arguments */
nir_store_var(&b, vars.accel_struct, nir_load_accel_struct_amd(&b), 1);
@@ -1587,9 +1567,7 @@ lower_rt_instruction_monolithic(nir_builder *b, nir_instr *instr, void *data)
case nir_intrinsic_execute_callable:
unreachable("nir_intrinsic_execute_callable");
case nir_intrinsic_trace_ray: {
nir_deref_instr *payload_deref = nir_src_as_deref(intr->src[10]);
assert(payload_deref->deref_type == nir_deref_type_var);
vars->arg = payload_deref->var;
nir_store_var(b, vars->arg, nir_iadd_imm(b, intr->src[10].ssa, -b->shader->scratch_size), 1);
nir_src cull_mask = intr->src[2];
bool ignore_cull_mask = nir_src_is_const(cull_mask) && (nir_src_as_uint(cull_mask) & 0xFF) == 0xFF;
@@ -1631,20 +1609,6 @@ lower_rt_instruction_monolithic(nir_builder *b, nir_instr *instr, void *data)
}
}
static bool
radv_rewrite_call_data_deref_modes(nir_builder *b, nir_instr *instr, void *data)
{
if (instr->type != nir_instr_type_deref)
return false;
nir_deref_instr *deref = nir_instr_as_deref(instr);
if (!nir_deref_mode_is(deref, nir_var_shader_call_data))
return false;
deref->modes = nir_var_shader_temp;
return true;
}
static void
lower_rt_instructions_monolithic(nir_shader *shader, struct radv_device *device,
struct radv_ray_tracing_pipeline *pipeline,
@@ -1662,8 +1626,6 @@ lower_rt_instructions_monolithic(nir_shader *shader, struct radv_device *device,
nir_shader_instructions_pass(shader, lower_rt_instruction_monolithic, nir_metadata_none, &state);
nir_index_ssa_defs(impl);
nir_shader_instructions_pass(shader, radv_rewrite_call_data_deref_modes, nir_metadata_all, &state);
/* Register storage for hit attributes */
nir_variable *hit_attribs[RADV_MAX_HIT_ATTRIB_SIZE / sizeof(uint32_t)];
@@ -1720,7 +1682,7 @@ radv_nir_lower_rt_abi(nir_shader *shader, const VkRayTracingPipelineCreateInfoKH
const VkPipelineCreateFlagBits2KHR create_flags = vk_rt_pipeline_create_flags(pCreateInfo);
struct rt_variables vars = create_rt_variables(shader, device, create_flags, monolithic);
struct rt_variables vars = create_rt_variables(shader, device, create_flags);
if (monolithic)
lower_rt_instructions_monolithic(shader, device, pipeline, pCreateInfo, &vars);
@@ -1751,9 +1713,7 @@ radv_nir_lower_rt_abi(nir_shader *shader, const VkRayTracingPipelineCreateInfoKH
nir_store_var(&b, vars.stack_ptr, ac_nir_load_arg(&b, &args->ac, args->ac.rt.dynamic_callable_stack_base), 1);
nir_def *record_ptr = ac_nir_load_arg(&b, &args->ac, args->ac.rt.shader_record);
nir_store_var(&b, vars.shader_record_ptr, nir_pack_64_2x32(&b, record_ptr), 1);
if (!monolithic)
nir_store_var(&b, vars.arg, ac_nir_load_arg(&b, &args->ac, args->ac.rt.payload_offset), 1);
nir_store_var(&b, vars.arg, ac_nir_load_arg(&b, &args->ac, args->ac.rt.payload_offset), 1);
nir_def *accel_struct = ac_nir_load_arg(&b, &args->ac, args->ac.rt.accel_struct);
nir_store_var(&b, vars.accel_struct, nir_pack_64_2x32(&b, accel_struct), 1);
@@ -1826,7 +1786,4 @@ radv_nir_lower_rt_abi(nir_shader *shader, const VkRayTracingPipelineCreateInfoKH
NIR_PASS_V(shader, nir_lower_vars_to_ssa);
if (shader->info.stage == MESA_SHADER_CLOSEST_HIT || shader->info.stage == MESA_SHADER_INTERSECTION)
NIR_PASS_V(shader, lower_hit_attribs, NULL, info->wave_size);
if (monolithic)
ac_nir_lower_indirect_derefs(shader, device->physical_device->rad_info.gfx_level);
}

7
src/amd/vulkan/radv_pipeline_rt.c
View File

@@ -24,7 +24,6 @@
#include "nir/nir.h"
#include "nir/nir_builder.h"
#include "nir/radv_nir.h"
#include "radv_debug.h"
#include "radv_private.h"
#include "radv_shader.h"
@@ -366,8 +365,6 @@ radv_rt_nir_to_asm(struct radv_device *device, struct vk_pipeline_cache *cache,
bool keep_executable_info = radv_pipeline_capture_shaders(device, pipeline->base.base.create_flags);
bool keep_statistic_info = radv_pipeline_capture_shader_stats(device, pipeline->base.base.create_flags);
radv_nir_lower_rt_io(stage->nir, monolithic);
/* Gather shader info. */
nir_shader_gather_info(stage->nir, nir_shader_get_entrypoint(stage->nir));
radv_nir_shader_info_init(stage->stage, MESA_SHADER_NONE, &stage->info);
@@ -520,9 +517,7 @@ radv_rt_compile_shaders(struct radv_device *device, struct vk_pipeline_cache *ca
radv_pipeline_stage_init(&pCreateInfo->pStages[i], pipeline_layout, stage);
/* precompile the shader */
stage->nir = radv_shader_spirv_to_nir(device, stage, key, false);
NIR_PASS(_, stage->nir, radv_nir_lower_hit_attrib_derefs);
stage->nir = radv_parse_rt_stage(device, &pCreateInfo->pStages[i], key, pipeline_layout);
rt_stages[i].can_inline = radv_rt_can_inline_shader(stage->nir);

4
src/amd/vulkan/radv_shader.h
View File

@@ -726,7 +726,9 @@ void radv_postprocess_nir(struct radv_device *device, const struct radv_pipeline
bool radv_shader_should_clear_lds(const struct radv_device *device, const nir_shader *shader);
void radv_nir_lower_rt_io(nir_shader *shader, bool monolithic);
nir_shader *radv_parse_rt_stage(struct radv_device *device, const VkPipelineShaderStageCreateInfo *sinfo,
const struct radv_pipeline_key *key,
const struct radv_pipeline_layout *pipeline_layout);
void radv_nir_lower_rt_abi(nir_shader *shader, const VkRayTracingPipelineCreateInfoKHR *pCreateInfo,
const struct radv_shader_args *args, const struct radv_shader_info *info,