radv: Use correct watermark for early loop exit.

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>

diff --git a/src/amd/vulkan/radv_nir_lower_ray_queries.c b/src/amd/vulkan/radv_nir_lower_ray_queries.c
index de2bc07..9de26a0 100644 (file)
--- a/src/amd/vulkan/radv_nir_lower_ray_queries.c
+++ b/src/amd/vulkan/radv_nir_lower_ray_queries.c
@@ -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));

diff --git a/src/amd/vulkan/radv_rt_common.c b/src/amd/vulkan/radv_rt_common.c
index d1eff47..6e1b74b 100644 (file)
--- a/src/amd/vulkan/radv_rt_common.c
+++ b/src/amd/vulkan/radv_rt_common.c
@@ -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);

diff --git a/src/amd/vulkan/radv_rt_common.h b/src/amd/vulkan/radv_rt_common.h
index 351fd27..1b748a1 100644 (file)
--- a/src/amd/vulkan/radv_rt_common.h
+++ b/src/amd/vulkan/radv_rt_common.h
@@ -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;

diff --git a/src/amd/vulkan/radv_rt_shader.c b/src/amd/vulkan/radv_rt_shader.c
index ab66657..52c5833 100644 (file)
--- a/src/amd/vulkan/radv_rt_shader.c
+++ b/src/amd/vulkan/radv_rt_shader.c
@@ -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)