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

radv: Use correct watermark for early loop exit.

Browse Source 
The previous check assumed the stack starts at offset=0, which isn't
necessarily true for ray queries.

Note that this didn't cause correctness issues, just made an optimization
not apply. Found when I accidentally made this load-bearing in a
refactor.

Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/20265>
This commit is contained in:
Bas Nieuwenhuizen
2022-12-11 03:08:55 +01:00
committed by Marge Bot
parent f0d6a1a685
commit efa4e9568b
4 changed files with 6 additions and 2 deletions
Download Patch File Download Diff File Expand all files Collapse all files

2
src/amd/vulkan/radv_nir_lower_ray_queries.c
View File

@@ -665,12 +665,14 @@ lower_rq_proceed(nir_builder *b, nir_ssa_def *index, struct ray_query_vars *vars
if (vars->stack) {
args.stack_stride = 1;
args.stack_entries = MAX_SCRATCH_STACK_ENTRY_COUNT;
args.stack_base = 0;
} else {
uint32_t workgroup_size = b->shader->info.workgroup_size[0] *
b->shader->info.workgroup_size[1] *
b->shader->info.workgroup_size[2];
args.stack_stride = workgroup_size * 4;
args.stack_entries = MAX_SHARED_STACK_ENTRY_COUNT;
args.stack_base = vars->shared_base;
}
nir_push_if(b, rq_load_var(b, index, vars->incomplete));

2
src/amd/vulkan/radv_rt_common.c
View File

@@ -551,7 +551,7 @@ radv_build_ray_traversal(struct radv_device *device, nir_builder *b,
/* Early exit if we never overflowed the stack, to avoid having to backtrack to
* the root for no reason. */
nir_push_if(b, nir_ilt(b, nir_load_deref(b, args->vars.stack),
nir_imm_int(b, args->stack_stride)));
nir_imm_int(b, args->stack_base + args->stack_stride)));
{
nir_store_var(b, incomplete, nir_imm_bool(b, false), 0x1);
nir_jump(b, nir_jump_break);

3
src/amd/vulkan/radv_rt_common.h
View File

@@ -147,9 +147,10 @@ struct radv_ray_traversal_args {
struct radv_ray_traversal_vars vars;
/* The increment/decrement used for radv_ray_traversal_vars::stack, and how many entries are
* available. */
* available. stack_base is the base address of the stack. */
uint32_t stack_stride;
uint32_t stack_entries;
uint32_t stack_base;
radv_rt_stack_store_cb stack_store_cb;
radv_rt_stack_load_cb stack_load_cb;

1
src/amd/vulkan/radv_rt_shader.c
View File

@@ -1371,6 +1371,7 @@ build_traversal_shader(struct radv_device *device,
.vars = trav_vars_args,
.stack_stride = device->physical_device->rt_wave_size * sizeof(uint32_t),
.stack_entries = MAX_STACK_ENTRY_COUNT,
.stack_base = 0,
.stack_store_cb = store_stack_entry,
.stack_load_cb = load_stack_entry,
.aabb_cb = (pCreateInfo->flags & VK_PIPELINE_CREATE_RAY_TRACING_SKIP_AABBS_BIT_KHR)