diff --git a/src/intel/compiler/elk/elk_compiler.h b/src/intel/compiler/elk/elk_compiler.h
index a3f26a847dc..128b3948267 100644
--- a/src/intel/compiler/elk/elk_compiler.h
+++ b/src/intel/compiler/elk/elk_compiler.h
@@ -754,9 +754,6 @@ struct elk_stage_prog_data {
    /** Does this program pull from any UBO or other constant buffers? */
    bool has_ubo_pull;
 
-   /** How many ray queries objects in this shader. */
-   unsigned ray_queries;
-
    /**
     * Register where the thread expects to find input data from the URB
     * (typically uniforms, followed by vertex or fragment attributes).
@@ -1209,7 +1206,6 @@ struct elk_cs_prog_data {
    bool uses_barrier;
    bool uses_num_work_groups;
    bool uses_inline_data;
-   bool uses_btd_stack_ids;
    uint8_t generate_local_id;
    enum intel_compute_walk_order walk_order;
 
diff --git a/src/intel/compiler/elk/elk_eu.h b/src/intel/compiler/elk/elk_eu.h
index 5c86775fe79..852da27322d 100644
--- a/src/intel/compiler/elk/elk_eu.h
+++ b/src/intel/compiler/elk/elk_eu.h
@@ -1674,18 +1674,6 @@ elk_btd_spawn_exec_size(UNUSED const struct intel_device_info *devinfo,
    return elk_mdc_sm2_exec_size(GET_BITS(desc, 8, 8));
 }
 
-static inline uint32_t
-elk_rt_trace_ray_desc(ASSERTED const struct intel_device_info *devinfo,
-                      unsigned exec_size)
-{
-   assert(devinfo->has_ray_tracing);
-   assert(devinfo->ver < 20 || exec_size == 16);
-
-   return SET_BITS(0, 19, 19) | /* No header */
-          SET_BITS(0, 17, 14) | /* Message type */
-          SET_BITS(elk_mdc_sm2(exec_size), 8, 8);
-}
-
 /**
  * Construct a message descriptor immediate with the specified pixel
  * interpolator function controls.
diff --git a/src/intel/compiler/elk/elk_eu_defines.h b/src/intel/compiler/elk/elk_eu_defines.h
index 8df64c1d424..904886af2a8 100644
--- a/src/intel/compiler/elk/elk_eu_defines.h
+++ b/src/intel/compiler/elk/elk_eu_defines.h
@@ -355,19 +355,6 @@ enum a64_logical_srcs {
    A64_LOGICAL_NUM_SRCS
 };
 
-enum rt_logical_srcs {
-   /** Address of the globals */
-   RT_LOGICAL_SRC_GLOBALS,
-   /** Level at which the tracing should start */
-   RT_LOGICAL_SRC_BVH_LEVEL,
-   /** Type of tracing operation */
-   RT_LOGICAL_SRC_TRACE_RAY_CONTROL,
-   /** Synchronous tracing (ray query) */
-   RT_LOGICAL_SRC_SYNCHRONOUS,
-
-   RT_LOGICAL_NUM_SRCS
-};
-
 enum urb_logical_srcs {
    URB_LOGICAL_SRC_HANDLE,
    URB_LOGICAL_SRC_PER_SLOT_OFFSETS,
diff --git a/src/intel/compiler/elk/elk_eu_opcodes.h b/src/intel/compiler/elk/elk_eu_opcodes.h
index 4993ddb6cd7..0f0622fef22 100644
--- a/src/intel/compiler/elk/elk_eu_opcodes.h
+++ b/src/intel/compiler/elk/elk_eu_opcodes.h
@@ -609,12 +609,7 @@ enum elk_opcode {
    ELK_TES_OPCODE_CREATE_INPUT_READ_HEADER,
    ELK_TES_OPCODE_ADD_INDIRECT_URB_OFFSET,
 
-   ELK_SHADER_OPCODE_BTD_SPAWN_LOGICAL,
-   ELK_SHADER_OPCODE_BTD_RETIRE_LOGICAL,
-
    ELK_SHADER_OPCODE_READ_SR_REG,
-
-   ELK_RT_OPCODE_TRACE_RAY_LOGICAL,
 };
 
 
diff --git a/src/intel/compiler/elk/elk_fs.cpp b/src/intel/compiler/elk/elk_fs.cpp
index 968540f46be..e4a29ba206c 100644
--- a/src/intel/compiler/elk/elk_fs.cpp
+++ b/src/intel/compiler/elk/elk_fs.cpp
@@ -7257,7 +7257,6 @@ elk_compile_fs(const struct elk_compiler *compiler,
                                    params->base.debug_flag : DEBUG_WM);
 
    prog_data->base.stage = MESA_SHADER_FRAGMENT;
-   prog_data->base.ray_queries = nir->info.ray_queries;
    prog_data->base.total_scratch = 0;
 
    const struct intel_device_info *devinfo = compiler->devinfo;
@@ -7689,7 +7688,6 @@ elk_compile_cs(const struct elk_compiler *compiler,
 
    prog_data->base.stage = MESA_SHADER_COMPUTE;
    prog_data->base.total_shared = nir->info.shared_size;
-   prog_data->base.ray_queries = nir->info.ray_queries;
    prog_data->base.total_scratch = 0;
 
    if (!nir->info.workgroup_size_variable) {
diff --git a/src/intel/compiler/elk/elk_fs_thread_payload.cpp b/src/intel/compiler/elk/elk_fs_thread_payload.cpp
index 58d52df0e80..90b6b6581fa 100644
--- a/src/intel/compiler/elk/elk_fs_thread_payload.cpp
+++ b/src/intel/compiler/elk/elk_fs_thread_payload.cpp
@@ -489,10 +489,6 @@ elk_cs_thread_payload::elk_cs_thread_payload(const elk_fs_visitor &v)
             local_invocation_id[i] = elk_imm_uw(0);
          }
       }
-
-      /* TODO: Fill out uses_btd_stack_ids automatically */
-      if (prog_data->uses_btd_stack_ids)
-         r += reg_unit(v.devinfo);
    }
 
    num_regs = r;
diff --git a/src/intel/compiler/elk/elk_nir.c b/src/intel/compiler/elk/elk_nir.c
index e0fd3f0da84..c320345b430 100644
--- a/src/intel/compiler/elk/elk_nir.c
+++ b/src/intel/compiler/elk/elk_nir.c
@@ -630,9 +630,6 @@ elk_nir_optimize(nir_shader *nir, bool is_scalar,
       OPT(nir_opt_dead_write_vars);
       OPT(nir_opt_combine_stores, nir_var_all);
 
-      OPT(nir_opt_ray_queries);
-      OPT(nir_opt_ray_query_ranges);
-
       if (is_scalar) {
          OPT(nir_lower_alu_to_scalar, NULL, NULL);
       } else {
diff --git a/src/intel/compiler/elk/elk_shader.cpp b/src/intel/compiler/elk/elk_shader.cpp
index 9cd16bccc91..7177e41c5ac 100644
--- a/src/intel/compiler/elk/elk_shader.cpp
+++ b/src/intel/compiler/elk/elk_shader.cpp
@@ -513,17 +513,10 @@ elk_instruction_name(const struct elk_isa_info *isa, enum elk_opcode op)
    case ELK_TES_OPCODE_GET_PRIMITIVE_ID:
       return "tes_get_primitive_id";
 
-   case ELK_RT_OPCODE_TRACE_RAY_LOGICAL:
-      return "rt_trace_ray_logical";
-
    case ELK_SHADER_OPCODE_RND_MODE:
       return "rnd_mode";
    case ELK_SHADER_OPCODE_FLOAT_CONTROL_MODE:
       return "float_control_mode";
-   case ELK_SHADER_OPCODE_BTD_SPAWN_LOGICAL:
-      return "btd_spawn_logical";
-   case ELK_SHADER_OPCODE_BTD_RETIRE_LOGICAL:
-      return "btd_retire_logical";
    case ELK_SHADER_OPCODE_READ_SR_REG:
       return "read_sr_reg";
    }
@@ -1102,9 +1095,6 @@ elk_backend_instruction::has_side_effects() const
    case ELK_FS_OPCODE_SCHEDULING_FENCE:
    case ELK_SHADER_OPCODE_OWORD_BLOCK_WRITE_LOGICAL:
    case ELK_SHADER_OPCODE_A64_OWORD_BLOCK_WRITE_LOGICAL:
-   case ELK_SHADER_OPCODE_BTD_SPAWN_LOGICAL:
-   case ELK_SHADER_OPCODE_BTD_RETIRE_LOGICAL:
-   case ELK_RT_OPCODE_TRACE_RAY_LOGICAL:
    case ELK_VEC4_OPCODE_ZERO_OOB_PUSH_REGS:
       return true;
    default:
@@ -1285,7 +1275,6 @@ elk_compile_tes(const struct elk_compiler *compiler,
    const unsigned *assembly;
 
    prog_data->base.base.stage = MESA_SHADER_TESS_EVAL;
-   prog_data->base.base.ray_queries = nir->info.ray_queries;
 
    nir->info.inputs_read = key->inputs_read;
    nir->info.patch_inputs_read = key->patch_inputs_read;
diff --git a/src/intel/compiler/elk/elk_simd_selection.cpp b/src/intel/compiler/elk/elk_simd_selection.cpp
index ad71f02a69a..08815e4a27c 100644
--- a/src/intel/compiler/elk/elk_simd_selection.cpp
+++ b/src/intel/compiler/elk/elk_simd_selection.cpp
@@ -132,16 +132,6 @@ elk_simd_should_compile(elk_simd_selection_state &state, unsigned simd)
       return false;
    }
 
-   if (width == 32 && cs_prog_data && cs_prog_data->base.ray_queries > 0) {
-      state.error[simd] = "Ray queries not supported";
-      return false;
-   }
-
-   if (width == 32 && cs_prog_data && cs_prog_data->uses_btd_stack_ids) {
-      state.error[simd] = "Bindless shader calls not supported";
-      return false;
-   }
-
    uint64_t start;
    switch (prog_data->stage) {
    case MESA_SHADER_COMPUTE:
diff --git a/src/intel/compiler/elk/elk_vec4.cpp b/src/intel/compiler/elk/elk_vec4.cpp
index f35335a7bf1..a51722e07e7 100644
--- a/src/intel/compiler/elk/elk_vec4.cpp
+++ b/src/intel/compiler/elk/elk_vec4.cpp
@@ -2559,7 +2559,6 @@ elk_compile_vs(const struct elk_compiler *compiler,
                                    params->base.debug_flag : DEBUG_VS);
 
    prog_data->base.base.stage = MESA_SHADER_VERTEX;
-   prog_data->base.base.ray_queries = nir->info.ray_queries;
    prog_data->base.base.total_scratch = 0;
 
    const bool is_scalar = compiler->scalar_stage[MESA_SHADER_VERTEX];
diff --git a/src/intel/compiler/elk/elk_vec4_gs_visitor.cpp b/src/intel/compiler/elk/elk_vec4_gs_visitor.cpp
index 7df5935d1e0..21867747b4e 100644
--- a/src/intel/compiler/elk/elk_vec4_gs_visitor.cpp
+++ b/src/intel/compiler/elk/elk_vec4_gs_visitor.cpp
@@ -597,7 +597,6 @@ elk_compile_gs(const struct elk_compiler *compiler,
    const bool debug_enabled = elk_should_print_shader(nir, DEBUG_GS);
 
    prog_data->base.base.stage = MESA_SHADER_GEOMETRY;
-   prog_data->base.base.ray_queries = nir->info.ray_queries;
    prog_data->base.base.total_scratch = 0;
 
    /* The GLSL linker will have already matched up GS inputs and the outputs
diff --git a/src/intel/compiler/elk/elk_vec4_tcs.cpp b/src/intel/compiler/elk/elk_vec4_tcs.cpp
index 919baab6516..be28602c0c6 100644
--- a/src/intel/compiler/elk/elk_vec4_tcs.cpp
+++ b/src/intel/compiler/elk/elk_vec4_tcs.cpp
@@ -366,7 +366,6 @@ elk_compile_tcs(const struct elk_compiler *compiler,
    const unsigned *assembly;
 
    vue_prog_data->base.stage = MESA_SHADER_TESS_CTRL;
-   prog_data->base.base.ray_queries = nir->info.ray_queries;
    prog_data->base.base.total_scratch = 0;
 
    nir->info.outputs_written = key->outputs_written;