intel/fs: fixup simd selection with shader calls

Signed-off-by: Lionel Landwerlin <lionel.g.landwerlin@intel.com>
Reviewed-by: Caio Oliveira <caio.oliveira@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/17908>

src/intel/compiler/brw_fs.cpp

@@ -7825,9 +7825,6 @@ brw_compile_cs(const struct brw_compiler *compiler,
                                   nir->info.workgroup_size_variable;
 
       if (v[simd]->run_cs(allow_spilling)) {
-         /* We should always be able to do SIMD32 for compute shaders. */
-         assert(v[simd]->max_dispatch_width >= 32);
-
          cs_fill_push_const_info(compiler->devinfo, prog_data);
 
          brw_simd_mark_compiled(simd, prog_data, v[simd]->spilled_any_registers);

src/intel/compiler/brw_simd_selection.c

@@ -60,8 +60,9 @@ brw_simd_should_compile(void *mem_ctx,
 
    const unsigned width = 8u << simd;
 
-   /* For shaders with variable size workgroup, we will always compile all the
+   /* For shaders with variable size workgroup, in most cases we can compile
-    * variants, since the choice will happen only at dispatch time.
+    * all the variants (exceptions are bindless dispatch & ray queries), since
+    * the choice will happen only at dispatch time.
     */
    const bool workgroup_size_variable = prog_data->local_size[0] == 0;
 
@@ -113,6 +114,20 @@ brw_simd_should_compile(void *mem_ctx,
       }
    }
 
+   if (width == 32 && prog_data->base.ray_queries > 0) {
+      *error = ralloc_asprintf(
+         mem_ctx, "SIMD%u skipped because of ray queries",
+         width);
+      return false;
+   }
+
+   if (width == 32 && prog_data->uses_btd_stack_ids) {
+      *error = ralloc_asprintf(
+         mem_ctx, "SIMD%u skipped because of bindless shader calls",
+         width);
+      return false;
+   }
+
    const bool env_skip[3] = {
       INTEL_DEBUG(DEBUG_NO8),
       INTEL_DEBUG(DEBUG_NO16),