diff --git a/src/amd/vulkan/radv_nir_lower_ray_queries.c b/src/amd/vulkan/radv_nir_lower_ray_queries.c
index de2bc07c0a5..9de26a054af 100644
--- 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 d1eff476392..6e1b74b787e 100644
--- 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 351fd27f5b4..1b748a1cc9f 100644
--- 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 ab666579daf..52c5833cd96 100644
--- 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)