Module: Mesa Branch: main Commit: 734384e8bc2606f8a69d281cdc4fe728070acab9 URL: http://cgit.freedesktop.org/mesa/mesa/commit/?id=734384e8bc2606f8a69d281cdc4fe728070acab9
Author: Lionel Landwerlin <[email protected]> Date: Wed Nov 3 12:42:29 2021 +0200 intel/fs: fixup simd selection with shader calls Signed-off-by: Lionel Landwerlin <[email protected]> Reviewed-by: Caio Oliveira <[email protected]> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/17908> --- src/intel/compiler/brw_fs.cpp | 3 --- src/intel/compiler/brw_simd_selection.c | 19 +++++++++++++++++-- 2 files changed, 17 insertions(+), 5 deletions(-) diff --git a/src/intel/compiler/brw_fs.cpp b/src/intel/compiler/brw_fs.cpp index 8010d62b9ef..0caa7413375 100644 --- a/src/intel/compiler/brw_fs.cpp +++ b/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); diff --git a/src/intel/compiler/brw_simd_selection.c b/src/intel/compiler/brw_simd_selection.c index 48055eb9b42..558f4ccc734 100644 --- a/src/intel/compiler/brw_simd_selection.c +++ b/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 - * variants, since the choice will happen only at dispatch time. + /* For shaders with variable size workgroup, in most cases we can compile + * 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),
