diff --git a/src/amd/compiler/aco_nir_call_attribs.h b/src/amd/compiler/aco_nir_call_attribs.h
index 777f1063c05..7c6572d1616 100644
--- a/src/amd/compiler/aco_nir_call_attribs.h
+++ b/src/amd/compiler/aco_nir_call_attribs.h
@@ -43,6 +43,8 @@ enum aco_nir_rt_function_arg {
    RT_ARG_LAUNCH_SIZE,
    RT_ARG_DESCRIPTORS,
    RT_ARG_DYNAMIC_DESCRIPTORS,
+   RT_ARG_HEAP_RESOURCE = RT_ARG_DESCRIPTORS,
+   RT_ARG_HEAP_SAMPLER = RT_ARG_DYNAMIC_DESCRIPTORS,
    RT_ARG_PUSH_CONSTANTS,
    RT_ARG_SBT_DESCRIPTORS,
    RT_ARG_COUNT,
diff --git a/src/amd/compiler/aco_shader_info.h b/src/amd/compiler/aco_shader_info.h
index 9467802645d..56f230b63b6 100644
--- a/src/amd/compiler/aco_shader_info.h
+++ b/src/amd/compiler/aco_shader_info.h
@@ -114,6 +114,7 @@ struct aco_shader_info {
    unsigned workgroup_size;
    unsigned lds_size;
    bool merged_shader_compiled_separately; /* GFX9+ */
+   bool descriptor_heap;
    struct ac_arg next_stage_pc;
    struct ac_arg epilog_pc; /* Vulkan only */
    struct {
diff --git a/src/amd/compiler/instruction_selection/aco_instruction_selection.h b/src/amd/compiler/instruction_selection/aco_instruction_selection.h
index 1caf41c6b49..a2390cd845a 100644
--- a/src/amd/compiler/instruction_selection/aco_instruction_selection.h
+++ b/src/amd/compiler/instruction_selection/aco_instruction_selection.h
@@ -306,7 +306,8 @@ void finish_program(isel_context* ctx);
 
 ABI nir_abi_to_aco(unsigned nir_abi_mask);
 
-param_assignment_hints get_ahit_isec_param_hints(const struct callee_info& traversal_info);
+param_assignment_hints get_ahit_isec_param_hints(const struct callee_info& traversal_info,
+                                                 bool uses_descriptor_heap);
 
 struct callee_info get_callee_info(amd_gfx_level gfx_level, unsigned wave_size, const ABI& abi,
                                    unsigned param_count, const nir_parameter* parameters,
diff --git a/src/amd/compiler/instruction_selection/aco_isel_helpers.cpp b/src/amd/compiler/instruction_selection/aco_isel_helpers.cpp
index d78f949af8a..7f62f98d524 100644
--- a/src/amd/compiler/instruction_selection/aco_isel_helpers.cpp
+++ b/src/amd/compiler/instruction_selection/aco_isel_helpers.cpp
@@ -902,7 +902,7 @@ param_hint_map(param_assignment_hints& hints, const struct callee_info& traversa
 }
 
 param_assignment_hints
-get_ahit_isec_param_hints(const struct callee_info& traversal_info)
+get_ahit_isec_param_hints(const struct callee_info& traversal_info, bool uses_descriptor_heap)
 {
    param_assignment_hints hints;
    hints.stack_pointer_affinity = traversal_info.stack_ptr;
@@ -914,8 +914,13 @@ get_ahit_isec_param_hints(const struct callee_info& traversal_info)
 
    param_hint_map(hints, traversal_info, RT_ARG_LAUNCH_ID, RT_ARG_LAUNCH_ID);
    param_hint_map(hints, traversal_info, RT_ARG_LAUNCH_SIZE, RT_ARG_LAUNCH_SIZE);
-   param_hint_map(hints, traversal_info, RT_ARG_DESCRIPTORS, RT_ARG_DESCRIPTORS);
-   param_hint_map(hints, traversal_info, RT_ARG_DYNAMIC_DESCRIPTORS, RT_ARG_DYNAMIC_DESCRIPTORS);
+   if (uses_descriptor_heap) {
+      param_hint_map(hints, traversal_info, RT_ARG_HEAP_RESOURCE, RT_ARG_HEAP_RESOURCE);
+      param_hint_map(hints, traversal_info, RT_ARG_HEAP_SAMPLER, RT_ARG_HEAP_SAMPLER);
+   } else {
+      param_hint_map(hints, traversal_info, RT_ARG_DESCRIPTORS, RT_ARG_DESCRIPTORS);
+      param_hint_map(hints, traversal_info, RT_ARG_DYNAMIC_DESCRIPTORS, RT_ARG_DYNAMIC_DESCRIPTORS);
+   }
    param_hint_map(hints, traversal_info, RT_ARG_PUSH_CONSTANTS, RT_ARG_PUSH_CONSTANTS);
    param_hint_map(hints, traversal_info, RT_ARG_SBT_DESCRIPTORS, RT_ARG_SBT_DESCRIPTORS);
    param_hint_map(hints, traversal_info, AHIT_ISEC_ARG_SHADER_RECORD_PTR,
diff --git a/src/amd/compiler/instruction_selection/aco_select_nir.cpp b/src/amd/compiler/instruction_selection/aco_select_nir.cpp
index dfddba4c971..b6ac174b17d 100644
--- a/src/amd/compiler/instruction_selection/aco_select_nir.cpp
+++ b/src/amd/compiler/instruction_selection/aco_select_nir.cpp
@@ -754,7 +754,7 @@ visit_call(isel_context* ctx, nir_call_instr* instr)
    param_assignment_hints hints;
 
    if (nir_abi == ACO_NIR_CALL_ABI_AHIT_ISEC)
-      hints = get_ahit_isec_param_hints(ctx->callee_info);
+      hints = get_ahit_isec_param_hints(ctx->callee_info, ctx->program->info.descriptor_heap);
 
    ABI abi = nir_abi_to_aco(instr->callee->driver_attributes);
 
@@ -1390,7 +1390,8 @@ select_program_rt(isel_context& ctx, unsigned shader_count, struct nir_shader* c
          callee_info traversal_info = get_callee_info(
             ctx.program->gfx_level, ctx.program->wave_size, rtTraversalABI,
             traversal_function->num_params, traversal_function->params, NULL, limit);
-         callee_hints = get_ahit_isec_param_hints(traversal_info);
+         callee_hints =
+            get_ahit_isec_param_hints(traversal_info, ctx.program->info.descriptor_heap);
       }
 
       /* TODO: callable abi? */
diff --git a/src/amd/vulkan/nir/radv_nir_rt_stage_common.c b/src/amd/vulkan/nir/radv_nir_rt_stage_common.c
index 80a412ad687..e913cd14739 100644
--- a/src/amd/vulkan/nir/radv_nir_rt_stage_common.c
+++ b/src/amd/vulkan/nir/radv_nir_rt_stage_common.c
@@ -257,7 +257,7 @@ radv_nir_return_param_from_type(nir_parameter *param, const glsl_type *type, boo
 }
 
 void
-radv_build_rt_prolog(struct radv_device *device, struct radv_shader_stage *stage)
+radv_build_rt_prolog(struct radv_device *device, struct radv_shader_stage *stage, bool uses_descriptor_heap)
 {
    const struct radv_physical_device *pdev = radv_device_physical(device);
 
@@ -268,6 +268,7 @@ radv_build_rt_prolog(struct radv_device *device, struct radv_shader_stage *stage
    stage->info.loads_push_constants = true;
    stage->info.loads_dynamic_offsets = true;
    stage->info.force_indirect_descriptors = true;
+   stage->info.descriptor_heap = uses_descriptor_heap;
    stage->info.wave_size = pdev->rt_wave_size;
    stage->info.workgroup_size = stage->info.wave_size;
    stage->info.user_data_0 = R_00B900_COMPUTE_USER_DATA_0;
@@ -288,11 +289,18 @@ radv_build_rt_prolog(struct radv_device *device, struct radv_shader_stage *stage
    b.shader->info.min_subgroup_size = pdev->rt_wave_size;
 
    nir_function *raygen_function = nir_function_create(b.shader, "raygen_func");
-   radv_nir_init_rt_function_params(raygen_function, MESA_SHADER_RAYGEN, 0, 0);
+   radv_nir_init_rt_function_params(raygen_function, MESA_SHADER_RAYGEN, 0, 0, uses_descriptor_heap);
+
+   nir_def *descriptors, *dynamic_descriptors, *heap_resource, *heap_sampler;
+   if (uses_descriptor_heap) {
+      heap_resource = ac_nir_load_arg(&b, &stage->args.ac, stage->args.descriptors[RADV_HEAP_RESOURCE]);
+      heap_sampler = ac_nir_load_arg(&b, &stage->args.ac, stage->args.descriptors[RADV_HEAP_SAMPLER]);
+   } else {
+      descriptors = ac_nir_load_arg(&b, &stage->args.ac, stage->args.descriptors[0]);
+      dynamic_descriptors = ac_nir_load_arg(&b, &stage->args.ac, stage->args.ac.dynamic_descriptors);
+   }
 
-   nir_def *descriptors = ac_nir_load_arg(&b, &stage->args.ac, stage->args.descriptors[0]);
    nir_def *push_constants = ac_nir_load_arg(&b, &stage->args.ac, stage->args.ac.push_constants);
-   nir_def *dynamic_descriptors = ac_nir_load_arg(&b, &stage->args.ac, stage->args.ac.dynamic_descriptors);
    nir_def *sbt_desc = nir_pack_64_2x32(&b, ac_nir_load_arg(&b, &stage->args.ac, stage->args.ac.rt.sbt_descriptors));
    nir_def *launch_size_addr = nir_pack_64_2x32(&b, ac_nir_load_arg(&b, &stage->args.ac, stage->args.ac.rt.launch_size_addr));
    nir_def *traversal_addr =
@@ -399,8 +407,13 @@ radv_build_rt_prolog(struct radv_device *device, struct radv_shader_stage *stage
    nir_def *params[RAYGEN_ARG_COUNT];
    params[RT_ARG_LAUNCH_ID] = nir_vec3(&b, id_x, id_y, wg_ids[2]);
    params[RT_ARG_LAUNCH_SIZE] = launch_sizes;
-   params[RT_ARG_DESCRIPTORS] = descriptors;
-   params[RT_ARG_DYNAMIC_DESCRIPTORS] = dynamic_descriptors;
+   if (uses_descriptor_heap) {
+      params[RT_ARG_HEAP_RESOURCE] = heap_resource;
+      params[RT_ARG_HEAP_SAMPLER] = heap_sampler;
+   } else {
+      params[RT_ARG_DESCRIPTORS] = descriptors;
+      params[RT_ARG_DYNAMIC_DESCRIPTORS] = dynamic_descriptors;
+   }
    params[RT_ARG_PUSH_CONSTANTS] = push_constants;
    params[RT_ARG_SBT_DESCRIPTORS] = sbt_desc;
    params[RAYGEN_ARG_SHADER_RECORD_PTR] = shader_record_ptr;
diff --git a/src/amd/vulkan/nir/radv_nir_rt_stage_common.h b/src/amd/vulkan/nir/radv_nir_rt_stage_common.h
index 0656b714bb0..2e01666ae0d 100644
--- a/src/amd/vulkan/nir/radv_nir_rt_stage_common.h
+++ b/src/amd/vulkan/nir/radv_nir_rt_stage_common.h
@@ -159,5 +159,6 @@ struct radv_nir_rt_traversal_result radv_build_traversal(struct radv_device *dev
                                                          struct radv_nir_rt_traversal_params *params,
                                                          struct radv_ray_tracing_stage_info *info);
 
-void radv_build_rt_prolog(struct radv_device *device, struct radv_shader_stage *stage);
+void radv_build_rt_prolog(struct radv_device *device, struct radv_shader_stage *stage, bool uses_descriptor_heap);
+
 #endif // MESA_RADV_NIR_RT_STAGE_COMMON_H
diff --git a/src/amd/vulkan/nir/radv_nir_rt_stage_cps.c b/src/amd/vulkan/nir/radv_nir_rt_stage_cps.c
index d554607fa91..27873502a30 100644
--- a/src/amd/vulkan/nir/radv_nir_rt_stage_cps.c
+++ b/src/amd/vulkan/nir/radv_nir_rt_stage_cps.c
@@ -507,15 +507,20 @@ radv_nir_lower_rt_io_cps(nir_shader *nir)
 }
 
 static void
-init_cps_function(nir_function *function, bool has_position_fetch)
+init_cps_function(nir_function *function, bool has_position_fetch, bool uses_descriptor_heap)
 {
    function->num_params = has_position_fetch ? CPS_ARG_COUNT : CPS_ARG_COUNT - 1;
    function->params = rzalloc_array_size(function->shader, sizeof(nir_parameter), function->num_params);
 
    radv_nir_param_from_type(function->params + RT_ARG_LAUNCH_ID, glsl_vector_type(GLSL_TYPE_UINT, 3), false, 0);
    radv_nir_param_from_type(function->params + RT_ARG_LAUNCH_SIZE, glsl_vector_type(GLSL_TYPE_UINT, 3), true, 0);
-   radv_nir_param_from_type(function->params + RT_ARG_DESCRIPTORS, glsl_uint_type(), true, 0);
-   radv_nir_param_from_type(function->params + RT_ARG_DYNAMIC_DESCRIPTORS, glsl_uint_type(), true, 0);
+   if (uses_descriptor_heap) {
+      radv_nir_param_from_type(function->params + RT_ARG_HEAP_RESOURCE, glsl_uint_type(), true, 0);
+      radv_nir_param_from_type(function->params + RT_ARG_HEAP_SAMPLER, glsl_uint_type(), true, 0);
+   } else {
+      radv_nir_param_from_type(function->params + RT_ARG_DESCRIPTORS, glsl_uint_type(), true, 0);
+      radv_nir_param_from_type(function->params + RT_ARG_DYNAMIC_DESCRIPTORS, glsl_uint_type(), true, 0);
+   }
    radv_nir_param_from_type(function->params + RT_ARG_PUSH_CONSTANTS, glsl_uint_type(), true, 0);
    radv_nir_param_from_type(function->params + RT_ARG_SBT_DESCRIPTORS, glsl_uint64_t_type(), true, 0);
    radv_nir_param_from_type(function->params + RAYGEN_ARG_TRAVERSAL_ADDR, glsl_uint64_t_type(), true, 0);
@@ -552,15 +557,16 @@ radv_nir_lower_rt_abi_cps(nir_shader *shader, const struct radv_shader_info *inf
                           struct radv_device *device, struct radv_ray_tracing_pipeline *pipeline,
                           bool has_position_fetch, const struct radv_ray_tracing_stage_info *traversal_info)
 {
+   const bool uses_descriptor_heap = pipeline->base.base.create_flags & VK_PIPELINE_CREATE_2_DESCRIPTOR_HEAP_BIT_EXT;
    nir_function_impl *impl = nir_shader_get_entrypoint(shader);
 
    /* The first raygen shader gets called by the prolog with the standard raygen signature. Only shaders called by the
     * first shader can use the CPS function signature.
     */
    if (shader->info.stage != MESA_SHADER_RAYGEN || resume_shader)
-      init_cps_function(impl->function, has_position_fetch);
+      init_cps_function(impl->function, has_position_fetch, uses_descriptor_heap);
    else
-      radv_nir_init_rt_function_params(impl->function, MESA_SHADER_RAYGEN, 0, 0);
+      radv_nir_init_rt_function_params(impl->function, MESA_SHADER_RAYGEN, 0, 0, uses_descriptor_heap);
 
    if (traversal_info) {
       unsigned idx;
@@ -624,14 +630,19 @@ radv_nir_lower_rt_abi_cps(nir_shader *shader, const struct radv_shader_info *inf
    /* tail-call next shader */
    nir_def *shader_addr = nir_load_var(&b, vars.shader_addr);
    nir_function *continuation_func = nir_function_create(shader, "continuation_func");
-   init_cps_function(continuation_func, has_position_fetch);
+   init_cps_function(continuation_func, has_position_fetch, uses_descriptor_heap);
 
    unsigned param_count = continuation_func->num_params;
    nir_def **next_args = rzalloc_array_size(b.shader, sizeof(nir_def *), param_count);
    next_args[RT_ARG_LAUNCH_ID] = nir_load_param(&b, RT_ARG_LAUNCH_ID);
    next_args[RT_ARG_LAUNCH_SIZE] = nir_load_param(&b, RT_ARG_LAUNCH_SIZE);
-   next_args[RT_ARG_DESCRIPTORS] = nir_load_param(&b, RT_ARG_DESCRIPTORS);
-   next_args[RT_ARG_DYNAMIC_DESCRIPTORS] = nir_load_param(&b, RT_ARG_DYNAMIC_DESCRIPTORS);
+   if (uses_descriptor_heap) {
+      next_args[RT_ARG_HEAP_RESOURCE] = nir_load_param(&b, RT_ARG_HEAP_RESOURCE);
+      next_args[RT_ARG_HEAP_SAMPLER] = nir_load_param(&b, RT_ARG_HEAP_SAMPLER);
+   } else {
+      next_args[RT_ARG_DESCRIPTORS] = nir_load_param(&b, RT_ARG_DESCRIPTORS);
+      next_args[RT_ARG_DYNAMIC_DESCRIPTORS] = nir_load_param(&b, RT_ARG_DYNAMIC_DESCRIPTORS);
+   }
    next_args[RT_ARG_PUSH_CONSTANTS] = nir_load_param(&b, RT_ARG_PUSH_CONSTANTS);
    next_args[RT_ARG_SBT_DESCRIPTORS] = nir_load_param(&b, RT_ARG_SBT_DESCRIPTORS);
    next_args[RAYGEN_ARG_TRAVERSAL_ADDR] = nir_load_var(&b, vars.traversal_addr);
diff --git a/src/amd/vulkan/nir/radv_nir_rt_stage_functions.c b/src/amd/vulkan/nir/radv_nir_rt_stage_functions.c
index 342f715dedf..eb319588b7d 100644
--- a/src/amd/vulkan/nir/radv_nir_rt_stage_functions.c
+++ b/src/amd/vulkan/nir/radv_nir_rt_stage_functions.c
@@ -22,22 +22,27 @@
 #include "vk_pipeline.h"
 
 static void
-radv_nir_init_common_rt_params(nir_function *function)
+radv_nir_init_common_rt_params(nir_function *function, bool uses_descriptor_heap)
 {
    radv_nir_param_from_type(function->params + RT_ARG_LAUNCH_ID, glsl_vector_type(GLSL_TYPE_UINT, 3), false, 0);
    radv_nir_param_from_type(function->params + RT_ARG_LAUNCH_SIZE, glsl_vector_type(GLSL_TYPE_UINT, 3), true, 0);
-   radv_nir_param_from_type(function->params + RT_ARG_DESCRIPTORS, glsl_uint_type(), true, 0);
-   radv_nir_param_from_type(function->params + RT_ARG_DYNAMIC_DESCRIPTORS, glsl_uint_type(), true, 0);
+   if (uses_descriptor_heap) {
+      radv_nir_param_from_type(function->params + RT_ARG_HEAP_RESOURCE, glsl_uint_type(), true, 0);
+      radv_nir_param_from_type(function->params + RT_ARG_HEAP_SAMPLER, glsl_uint_type(), true, 0);
+   } else {
+      radv_nir_param_from_type(function->params + RT_ARG_DESCRIPTORS, glsl_uint_type(), true, 0);
+      radv_nir_param_from_type(function->params + RT_ARG_DYNAMIC_DESCRIPTORS, glsl_uint_type(), true, 0);
+   }
    radv_nir_param_from_type(function->params + RT_ARG_PUSH_CONSTANTS, glsl_uint_type(), true, 0);
    radv_nir_param_from_type(function->params + RT_ARG_SBT_DESCRIPTORS, glsl_uint64_t_type(), true, 0);
 }
 
 static void
-radv_nir_init_traversal_params(nir_function *function, unsigned payload_size)
+radv_nir_init_traversal_params(nir_function *function, unsigned payload_size, bool uses_descriptor_heap)
 {
    function->num_params = TRAVERSAL_ARG_PAYLOAD_BASE + DIV_ROUND_UP(payload_size, 4);
    function->params = rzalloc_array_size(function->shader, sizeof(nir_parameter), function->num_params);
-   radv_nir_init_common_rt_params(function);
+   radv_nir_init_common_rt_params(function, uses_descriptor_heap);
    radv_nir_param_from_type(function->params + TRAVERSAL_ARG_TRAVERSAL_ADDR, glsl_uint64_t_type(), true, 0);
    radv_nir_param_from_type(function->params + TRAVERSAL_ARG_SHADER_RECORD_PTR, glsl_uint64_t_type(), false, ACO_NIR_PARAM_ATTRIB_DISCARDABLE);
    radv_nir_param_from_type(function->params + TRAVERSAL_ARG_ACCEL_STRUCT, glsl_uint64_t_type(), false, 0);
@@ -68,7 +73,7 @@ radv_nir_init_traversal_params(nir_function *function, unsigned payload_size)
 
 void
 radv_nir_init_rt_function_params(nir_function *function, mesa_shader_stage stage, unsigned payload_size,
-                                 unsigned hit_attrib_size)
+                                 unsigned hit_attrib_size, bool uses_descriptor_heap)
 {
    unsigned payload_base = -1u;
 
@@ -76,7 +81,7 @@ radv_nir_init_rt_function_params(nir_function *function, mesa_shader_stage stage
    case MESA_SHADER_RAYGEN:
       function->num_params = RAYGEN_ARG_COUNT;
       function->params = rzalloc_array_size(function->shader, sizeof(nir_parameter), function->num_params);
-      radv_nir_init_common_rt_params(function);
+      radv_nir_init_common_rt_params(function, uses_descriptor_heap);
       radv_nir_param_from_type(function->params + RAYGEN_ARG_TRAVERSAL_ADDR, glsl_uint64_t_type(), true, 0);
       radv_nir_param_from_type(function->params + RAYGEN_ARG_SHADER_RECORD_PTR, glsl_uint64_t_type(), false, 0);
       function->driver_attributes = (uint32_t)ACO_NIR_CALL_ABI_RT_RECURSIVE | ACO_NIR_FUNCTION_ATTRIB_NORETURN;
@@ -84,7 +89,7 @@ radv_nir_init_rt_function_params(nir_function *function, mesa_shader_stage stage
    case MESA_SHADER_CALLABLE:
       function->num_params = RAYGEN_ARG_COUNT + DIV_ROUND_UP(payload_size, 4);
       function->params = rzalloc_array_size(function->shader, sizeof(nir_parameter), function->num_params);
-      radv_nir_init_common_rt_params(function);
+      radv_nir_init_common_rt_params(function, uses_descriptor_heap);
       radv_nir_param_from_type(function->params + RAYGEN_ARG_TRAVERSAL_ADDR, glsl_uint64_t_type(), true, 0);
       radv_nir_param_from_type(function->params + RAYGEN_ARG_SHADER_RECORD_PTR, glsl_uint64_t_type(), false, 0);
 
@@ -96,7 +101,7 @@ radv_nir_init_rt_function_params(nir_function *function, mesa_shader_stage stage
       function->num_params =
          AHIT_ISEC_ARG_HIT_ATTRIB_PAYLOAD_BASE + DIV_ROUND_UP(hit_attrib_size, 4) + DIV_ROUND_UP(payload_size, 4);
       function->params = rzalloc_array_size(function->shader, sizeof(nir_parameter), function->num_params);
-      radv_nir_init_common_rt_params(function);
+      radv_nir_init_common_rt_params(function, uses_descriptor_heap);
       radv_nir_param_from_type(function->params + AHIT_ISEC_ARG_SHADER_RECORD_PTR, glsl_uint64_t_type(), false, 0);
       radv_nir_param_from_type(function->params + AHIT_ISEC_ARG_CULL_MASK_AND_FLAGS, glsl_uint_type(), false, 0);
       radv_nir_param_from_type(function->params + AHIT_ISEC_ARG_SBT_INDEX, glsl_uint_type(), false, 0);
@@ -126,7 +131,7 @@ radv_nir_init_rt_function_params(nir_function *function, mesa_shader_stage stage
    case MESA_SHADER_MISS:
       function->num_params = CHIT_MISS_ARG_PAYLOAD_BASE + DIV_ROUND_UP(payload_size, 4);
       function->params = rzalloc_array_size(function->shader, sizeof(nir_parameter), function->num_params);
-      radv_nir_init_common_rt_params(function);
+      radv_nir_init_common_rt_params(function, uses_descriptor_heap);
       radv_nir_param_from_type(function->params + CHIT_MISS_ARG_TRAVERSAL_ADDR, glsl_uint64_t_type(), true, 0);
       radv_nir_param_from_type(function->params + CHIT_MISS_ARG_SHADER_RECORD_PTR, glsl_uint64_t_type(), false, 0);
       radv_nir_param_from_type(function->params + CHIT_MISS_ARG_ACCEL_STRUCT, glsl_uint64_t_type(), false, 0);
@@ -212,6 +217,7 @@ static struct rt_variables
 create_rt_variables(nir_shader *shader, struct radv_device *device, const VkPipelineCreateFlags2 flags,
                     unsigned max_payload_size, unsigned max_hit_attrib_size)
 {
+   const bool uses_descriptor_heap = flags & VK_PIPELINE_CREATE_2_DESCRIPTOR_HEAP_BIT_EXT;
    struct rt_variables vars = {
       .device = device,
       .flags = flags,
@@ -227,16 +233,19 @@ create_rt_variables(nir_shader *shader, struct radv_device *device, const VkPipe
    }
 
    nir_function *trace_ray_func = nir_function_create(shader, "trace_ray_func");
-   radv_nir_init_traversal_params(trace_ray_func, max_payload_size);
+   radv_nir_init_traversal_params(trace_ray_func, max_payload_size, uses_descriptor_heap);
    vars.trace_ray_func = trace_ray_func;
    nir_function *ahit_isec_func = nir_function_create(shader, "ahit_isec_func");
-   radv_nir_init_rt_function_params(ahit_isec_func, MESA_SHADER_ANY_HIT, max_payload_size, max_hit_attrib_size);
+   radv_nir_init_rt_function_params(ahit_isec_func, MESA_SHADER_ANY_HIT, max_payload_size, max_hit_attrib_size,
+                                    uses_descriptor_heap);
    vars.ahit_isec_func = ahit_isec_func;
    nir_function *chit_miss_func = nir_function_create(shader, "chit_miss_func");
-   radv_nir_init_rt_function_params(chit_miss_func, MESA_SHADER_CLOSEST_HIT, max_payload_size, max_hit_attrib_size);
+   radv_nir_init_rt_function_params(chit_miss_func, MESA_SHADER_CLOSEST_HIT, max_payload_size, max_hit_attrib_size,
+                                    uses_descriptor_heap);
    vars.chit_miss_func = chit_miss_func;
    nir_function *callable_func = nir_function_create(shader, "callable_func");
-   radv_nir_init_rt_function_params(callable_func, MESA_SHADER_CALLABLE, max_payload_size, max_hit_attrib_size);
+   radv_nir_init_rt_function_params(callable_func, MESA_SHADER_CALLABLE, max_payload_size, max_hit_attrib_size,
+                                    uses_descriptor_heap);
    vars.callable_func = callable_func;
 
    vars.shader_record_ptr_param = -1u;
@@ -339,6 +348,7 @@ lower_rt_instruction(nir_builder *b, nir_instr *instr, void *_vars)
    nir_intrinsic_instr *intr = nir_instr_as_intrinsic(instr);
 
    struct rt_variables *vars = _vars;
+   const bool uses_descriptor_heap = vars->flags & VK_PIPELINE_CREATE_2_DESCRIPTOR_HEAP_BIT_EXT;
 
    b->cursor = nir_before_instr(&intr->instr);
 
@@ -352,8 +362,13 @@ lower_rt_instruction(nir_builder *b, nir_instr *instr, void *_vars)
       nir_def **args = rzalloc_array_size(b->shader, sizeof(nir_def *), param_count);
       args[RT_ARG_LAUNCH_ID] = nir_load_param(b, RT_ARG_LAUNCH_ID);
       args[RT_ARG_LAUNCH_SIZE] = nir_load_param(b, RT_ARG_LAUNCH_SIZE);
-      args[RT_ARG_DESCRIPTORS] = nir_load_param(b, RT_ARG_DESCRIPTORS);
-      args[RT_ARG_DYNAMIC_DESCRIPTORS] = nir_load_param(b, RT_ARG_DYNAMIC_DESCRIPTORS);
+      if (uses_descriptor_heap) {
+         args[RT_ARG_HEAP_RESOURCE] = nir_load_param(b, RT_ARG_HEAP_RESOURCE);
+         args[RT_ARG_HEAP_SAMPLER] = nir_load_param(b, RT_ARG_HEAP_SAMPLER);
+      } else {
+         args[RT_ARG_DESCRIPTORS] = nir_load_param(b, RT_ARG_DESCRIPTORS);
+         args[RT_ARG_DYNAMIC_DESCRIPTORS] = nir_load_param(b, RT_ARG_DYNAMIC_DESCRIPTORS);
+      }
       args[RT_ARG_PUSH_CONSTANTS] = nir_load_param(b, RT_ARG_PUSH_CONSTANTS);
       args[RT_ARG_SBT_DESCRIPTORS] = nir_load_param(b, RT_ARG_SBT_DESCRIPTORS);
       args[RAYGEN_ARG_TRAVERSAL_ADDR] = nir_undef(b, 1, 64);
@@ -375,7 +390,13 @@ lower_rt_instruction(nir_builder *b, nir_instr *instr, void *_vars)
       args[RT_ARG_LAUNCH_ID] = nir_load_param(b, RT_ARG_LAUNCH_ID);
       args[RT_ARG_LAUNCH_SIZE] = nir_load_param(b, RT_ARG_LAUNCH_SIZE);
       args[RT_ARG_DESCRIPTORS] = nir_load_param(b, RT_ARG_DESCRIPTORS);
-      args[RT_ARG_DYNAMIC_DESCRIPTORS] = nir_load_param(b, RT_ARG_DYNAMIC_DESCRIPTORS);
+      if (uses_descriptor_heap) {
+         args[RT_ARG_HEAP_RESOURCE] = nir_load_param(b, RT_ARG_HEAP_RESOURCE);
+         args[RT_ARG_HEAP_SAMPLER] = nir_load_param(b, RT_ARG_HEAP_SAMPLER);
+      } else {
+         args[RT_ARG_DESCRIPTORS] = nir_load_param(b, RT_ARG_DESCRIPTORS);
+         args[RT_ARG_DYNAMIC_DESCRIPTORS] = nir_load_param(b, RT_ARG_DYNAMIC_DESCRIPTORS);
+      }
       args[RT_ARG_PUSH_CONSTANTS] = nir_load_param(b, RT_ARG_PUSH_CONSTANTS);
       args[RT_ARG_SBT_DESCRIPTORS] = nir_load_param(b, RT_ARG_SBT_DESCRIPTORS);
       args[TRAVERSAL_ARG_TRAVERSAL_ADDR] = traversal_addr;
@@ -503,6 +524,14 @@ lower_rt_instruction(nir_builder *b, nir_instr *instr, void *_vars)
       ret = nir_load_param(b, RT_ARG_DESCRIPTORS);
       break;
    }
+   case nir_intrinsic_load_rt_heap_resource_amd: {
+      ret = nir_load_param(b, RT_ARG_HEAP_RESOURCE);
+      break;
+   }
+   case nir_intrinsic_load_rt_heap_sampler_amd: {
+      ret = nir_load_param(b, RT_ARG_HEAP_SAMPLER);
+      break;
+   }
    case nir_intrinsic_load_rt_dynamic_descriptors_amd: {
       ret = nir_load_param(b, RT_ARG_DYNAMIC_DESCRIPTORS);
       break;
@@ -553,8 +582,13 @@ lower_rt_instruction(nir_builder *b, nir_instr *instr, void *_vars)
       nir_def **args = rzalloc_array_size(b->shader, sizeof(nir_def *), param_count);
       args[RT_ARG_LAUNCH_ID] = nir_load_param(b, RT_ARG_LAUNCH_ID);
       args[RT_ARG_LAUNCH_SIZE] = nir_load_param(b, RT_ARG_LAUNCH_SIZE);
-      args[RT_ARG_DESCRIPTORS] = nir_load_param(b, RT_ARG_DESCRIPTORS);
-      args[RT_ARG_DYNAMIC_DESCRIPTORS] = nir_load_param(b, RT_ARG_DYNAMIC_DESCRIPTORS);
+      if (uses_descriptor_heap) {
+         args[RT_ARG_HEAP_RESOURCE] = nir_load_param(b, RT_ARG_HEAP_RESOURCE);
+         args[RT_ARG_HEAP_SAMPLER] = nir_load_param(b, RT_ARG_HEAP_SAMPLER);
+      } else {
+         args[RT_ARG_DESCRIPTORS] = nir_load_param(b, RT_ARG_DESCRIPTORS);
+         args[RT_ARG_DYNAMIC_DESCRIPTORS] = nir_load_param(b, RT_ARG_DYNAMIC_DESCRIPTORS);
+      }
       args[RT_ARG_PUSH_CONSTANTS] = nir_load_param(b, RT_ARG_PUSH_CONSTANTS);
       args[RT_ARG_SBT_DESCRIPTORS] = nir_load_param(b, RT_ARG_SBT_DESCRIPTORS);
       args[CHIT_MISS_ARG_TRAVERSAL_ADDR] = nir_load_param(b, vars->traversal_addr_param);
@@ -598,8 +632,13 @@ lower_rt_instruction(nir_builder *b, nir_instr *instr, void *_vars)
       nir_def **args = rzalloc_array_size(b->shader, sizeof(nir_def *), param_count);
       args[RT_ARG_LAUNCH_ID] = nir_load_param(b, RT_ARG_LAUNCH_ID);
       args[RT_ARG_LAUNCH_SIZE] = nir_load_param(b, RT_ARG_LAUNCH_SIZE);
-      args[RT_ARG_DESCRIPTORS] = nir_load_param(b, RT_ARG_DESCRIPTORS);
-      args[RT_ARG_DYNAMIC_DESCRIPTORS] = nir_load_param(b, RT_ARG_DYNAMIC_DESCRIPTORS);
+      if (uses_descriptor_heap) {
+         args[RT_ARG_HEAP_RESOURCE] = nir_load_param(b, RT_ARG_HEAP_RESOURCE);
+         args[RT_ARG_HEAP_SAMPLER] = nir_load_param(b, RT_ARG_HEAP_SAMPLER);
+      } else {
+         args[RT_ARG_DESCRIPTORS] = nir_load_param(b, RT_ARG_DESCRIPTORS);
+         args[RT_ARG_DYNAMIC_DESCRIPTORS] = nir_load_param(b, RT_ARG_DYNAMIC_DESCRIPTORS);
+      }
       args[RT_ARG_PUSH_CONSTANTS] = nir_load_param(b, RT_ARG_PUSH_CONSTANTS);
       args[RT_ARG_SBT_DESCRIPTORS] = nir_load_param(b, RT_ARG_SBT_DESCRIPTORS);
       args[CHIT_MISS_ARG_TRAVERSAL_ADDR] = nir_load_param(b, vars->traversal_addr_param);
@@ -816,13 +855,15 @@ radv_nir_lower_rt_abi_functions(nir_shader *shader, const struct radv_shader_inf
                                 uint32_t hit_attrib_size, struct radv_device *device,
                                 struct radv_ray_tracing_pipeline *pipeline)
 {
+   const bool uses_descriptor_heap = pipeline->base.base.create_flags & VK_PIPELINE_CREATE_2_DESCRIPTOR_HEAP_BIT_EXT;
    nir_function_impl *impl = nir_shader_get_entrypoint(shader);
    nir_function *entrypoint_function = impl->function;
 
    if (radv_is_traversal_shader(shader))
-      radv_nir_init_traversal_params(entrypoint_function, payload_size);
+      radv_nir_init_traversal_params(entrypoint_function, payload_size, uses_descriptor_heap);
    else
-      radv_nir_init_rt_function_params(entrypoint_function, shader->info.stage, payload_size, hit_attrib_size);
+      radv_nir_init_rt_function_params(entrypoint_function, shader->info.stage, payload_size, hit_attrib_size,
+                                       uses_descriptor_heap);
 
    struct rt_variables vars =
       create_rt_variables(shader, device, pipeline->base.base.create_flags, payload_size, hit_attrib_size);
diff --git a/src/amd/vulkan/nir/radv_nir_rt_stage_functions.h b/src/amd/vulkan/nir/radv_nir_rt_stage_functions.h
index 0e19a3d5a27..b1f7d881832 100644
--- a/src/amd/vulkan/nir/radv_nir_rt_stage_functions.h
+++ b/src/amd/vulkan/nir/radv_nir_rt_stage_functions.h
@@ -14,7 +14,7 @@
 nir_function_impl *radv_get_rt_shader_entrypoint(nir_shader *shader);
 
 void radv_nir_init_rt_function_params(nir_function *function, mesa_shader_stage stage, unsigned payload_size,
-                                      unsigned hit_attrib_size);
+                                      unsigned hit_attrib_size, bool uses_descriptor_heap);
 
 void radv_nir_lower_rt_abi_functions(nir_shader *shader, const struct radv_shader_info *info, uint32_t payload_size,
                                      uint32_t hit_attrib_size, struct radv_device *device,
diff --git a/src/amd/vulkan/nir/radv_nir_rt_stage_monolithic.c b/src/amd/vulkan/nir/radv_nir_rt_stage_monolithic.c
index 3b966ed6357..2eefb844b4f 100644
--- a/src/amd/vulkan/nir/radv_nir_rt_stage_monolithic.c
+++ b/src/amd/vulkan/nir/radv_nir_rt_stage_monolithic.c
@@ -458,8 +458,9 @@ void
 radv_nir_lower_rt_abi_monolithic(nir_shader *shader, struct radv_device *device,
                                  struct radv_ray_tracing_pipeline *pipeline)
 {
+   const bool uses_descriptor_heap = pipeline->base.base.create_flags & VK_PIPELINE_CREATE_2_DESCRIPTOR_HEAP_BIT_EXT;
    nir_function_impl *impl = nir_shader_get_entrypoint(shader);
-   radv_nir_init_rt_function_params(impl->function, MESA_SHADER_RAYGEN, 0, 0);
+   radv_nir_init_rt_function_params(impl->function, MESA_SHADER_RAYGEN, 0, 0, uses_descriptor_heap);
 
    nir_builder b = nir_builder_at(nir_before_impl(impl));
 
diff --git a/src/amd/vulkan/nir/radv_nir_rt_traversal_shader.c b/src/amd/vulkan/nir/radv_nir_rt_traversal_shader.c
index dfe34d84cb9..1debc5c1649 100644
--- a/src/amd/vulkan/nir/radv_nir_rt_traversal_shader.c
+++ b/src/amd/vulkan/nir/radv_nir_rt_traversal_shader.c
@@ -834,6 +834,8 @@ handle_candidate_triangle(nir_builder *b, struct radv_triangle_intersection *int
                           const struct radv_ray_traversal_args *args, const struct radv_ray_flags *ray_flags)
 {
    struct traversal_data *data = args->data;
+   const bool uses_descriptor_heap =
+      data->pipeline->base.base.create_flags & VK_PIPELINE_CREATE_2_DESCRIPTOR_HEAP_BIT_EXT;
 
    nir_def *geometry_id = nir_iand_imm(b, intersection->base.geometry_id_and_flags, 0xfffffff);
    nir_def *sbt_idx =
@@ -925,8 +927,13 @@ handle_candidate_triangle(nir_builder *b, struct radv_triangle_intersection *int
          nir_def **params = rzalloc_array_size(b->shader, sizeof(nir_def *), param_count);
          params[RT_ARG_LAUNCH_ID] = nir_load_ray_launch_id(b);
          params[RT_ARG_LAUNCH_SIZE] = nir_load_ray_launch_size(b);
-         params[RT_ARG_DESCRIPTORS] = nir_load_rt_descriptors_amd(b);
-         params[RT_ARG_DYNAMIC_DESCRIPTORS] = nir_load_rt_dynamic_descriptors_amd(b);
+         if (uses_descriptor_heap) {
+            params[RT_ARG_HEAP_RESOURCE] = nir_load_rt_heap_resource_amd(b);
+            params[RT_ARG_HEAP_SAMPLER] = nir_load_rt_heap_sampler_amd(b);
+         } else {
+            params[RT_ARG_DESCRIPTORS] = nir_load_rt_descriptors_amd(b);
+            params[RT_ARG_DYNAMIC_DESCRIPTORS] = nir_load_rt_dynamic_descriptors_amd(b);
+         }
          params[RT_ARG_PUSH_CONSTANTS] = nir_load_rt_push_constants_amd(b);
          params[RT_ARG_SBT_DESCRIPTORS] = nir_load_sbt_base_amd(b);
          params[AHIT_ISEC_ARG_SHADER_RECORD_PTR] = sbt_data.shader_record_ptr;
@@ -988,6 +995,8 @@ handle_candidate_aabb(nir_builder *b, struct radv_leaf_intersection *intersectio
                       const struct radv_ray_traversal_args *args)
 {
    struct traversal_data *data = args->data;
+   const bool uses_descriptor_heap =
+      data->pipeline->base.base.create_flags & VK_PIPELINE_CREATE_2_DESCRIPTOR_HEAP_BIT_EXT;
 
    nir_def *geometry_id = nir_iand_imm(b, intersection->geometry_id_and_flags, 0xfffffff);
    nir_def *sbt_idx =
@@ -1075,8 +1084,13 @@ handle_candidate_aabb(nir_builder *b, struct radv_leaf_intersection *intersectio
       nir_def **params = rzalloc_array_size(b->shader, sizeof(nir_def *), param_count);
       params[RT_ARG_LAUNCH_ID] = nir_load_ray_launch_id(b);
       params[RT_ARG_LAUNCH_SIZE] = nir_load_ray_launch_size(b);
-      params[RT_ARG_DESCRIPTORS] = nir_load_rt_descriptors_amd(b);
-      params[RT_ARG_DYNAMIC_DESCRIPTORS] = nir_load_rt_dynamic_descriptors_amd(b);
+      if (uses_descriptor_heap) {
+         params[RT_ARG_HEAP_RESOURCE] = nir_load_rt_heap_resource_amd(b);
+         params[RT_ARG_HEAP_SAMPLER] = nir_load_rt_heap_sampler_amd(b);
+      } else {
+         params[RT_ARG_DESCRIPTORS] = nir_load_rt_descriptors_amd(b);
+         params[RT_ARG_DYNAMIC_DESCRIPTORS] = nir_load_rt_dynamic_descriptors_amd(b);
+      }
       params[RT_ARG_PUSH_CONSTANTS] = nir_load_rt_push_constants_amd(b);
       params[RT_ARG_SBT_DESCRIPTORS] = nir_load_sbt_base_amd(b);
       params[AHIT_ISEC_ARG_SHADER_RECORD_PTR] = sbt_data.shader_record_ptr;
@@ -1137,6 +1151,7 @@ struct radv_nir_rt_traversal_result
 radv_build_traversal(struct radv_device *device, struct radv_ray_tracing_pipeline *pipeline, nir_builder *b,
                      struct radv_nir_rt_traversal_params *params, struct radv_ray_tracing_stage_info *info)
 {
+   const bool uses_descriptor_heap = pipeline->base.base.create_flags & VK_PIPELINE_CREATE_2_DESCRIPTOR_HEAP_BIT_EXT;
    const struct radv_physical_device *pdev = radv_device_physical(device);
    nir_variable *barycentrics =
       nir_variable_create(b->shader, nir_var_ray_hit_attrib, glsl_vector_type(GLSL_TYPE_FLOAT, 2), "barycentrics");
@@ -1153,7 +1168,7 @@ radv_build_traversal(struct radv_device *device, struct radv_ray_tracing_pipelin
    if (!params->preprocess_ahit_isec) {
       nir_function *ahit_isec_func = nir_function_create(b->shader, "ahit_isec_func");
       radv_nir_init_rt_function_params(ahit_isec_func, MESA_SHADER_ANY_HIT, params->payload_size,
-                                       params->hit_attrib_size);
+                                       params->hit_attrib_size, uses_descriptor_heap);
       data.ahit_isec_func = ahit_isec_func;
    }
 
diff --git a/src/amd/vulkan/radv_aco_shader_info.h b/src/amd/vulkan/radv_aco_shader_info.h
index 5655d315b46..6c486284b2c 100644
--- a/src/amd/vulkan/radv_aco_shader_info.h
+++ b/src/amd/vulkan/radv_aco_shader_info.h
@@ -57,6 +57,7 @@ radv_aco_convert_shader_info(struct aco_shader_info *aco_info, const struct radv
    ASSIGN_FIELD(vs.has_prolog);
    ASSIGN_FIELD(ps.num_inputs);
    ASSIGN_FIELD(cs.uses_full_subgroups);
+   ASSIGN_FIELD(descriptor_heap);
    aco_info->vs.any_tcs_inputs_via_lds = radv->vs.tcs_inputs_via_lds != 0;
    /* S2 must not be modified for correct hang recovery when NGG_WAVE_ID_EN=1. */
    aco_info->vs.preserve_s2 = ngg_wave_id_en && gfx_level < GFX12;
diff --git a/src/amd/vulkan/radv_pipeline_rt.c b/src/amd/vulkan/radv_pipeline_rt.c
index 17366eb330f..c1de5a42a34 100644
--- a/src/amd/vulkan/radv_pipeline_rt.c
+++ b/src/amd/vulkan/radv_pipeline_rt.c
@@ -1044,11 +1044,12 @@ postprocess_rt_config(struct ac_shader_config *config, const struct radeon_info
 static void
 compile_rt_prolog(struct radv_device *device, struct radv_ray_tracing_pipeline *pipeline)
 {
+   const bool uses_descriptor_heap = pipeline->base.base.create_flags & VK_PIPELINE_CREATE_2_DESCRIPTOR_HEAP_BIT_EXT;
    const struct radv_physical_device *pdev = radv_device_physical(device);
    uint32_t push_constant_size = 0;
 
    struct radv_shader_stage prolog_stage = {0};
-   radv_build_rt_prolog(device, &prolog_stage);
+   radv_build_rt_prolog(device, &prolog_stage, uses_descriptor_heap);
    prolog_stage.nir->options = &pdev->nir_options[MESA_SHADER_COMPUTE];
    radv_optimize_nir(prolog_stage.nir, false);
    radv_postprocess_nir(device, NULL, &prolog_stage);