From 3afd572d8fd505b6d6b69a6ce08b7e523faa5f8f Mon Sep 17 00:00:00 2001 From: Alyssa Rosenzweig Date: Tue, 7 Apr 2026 14:20:45 -0400 Subject: [PATCH] brw: add Jay-specific SIMD selection rule In the future this might even do something clever. Signed-off-by: Alyssa Rosenzweig Reviewed-by: Lionel Landwerlin Reviewed-by: Kenneth Graunke Part-of: --- src/intel/compiler/brw/brw_shader.cpp | 13 ++++++++++++- 1 file changed, 12 insertions(+), 1 deletion(-) diff --git a/src/intel/compiler/brw/brw_shader.cpp b/src/intel/compiler/brw/brw_shader.cpp index e0f91b81dba..3f85615413d 100644 --- a/src/intel/compiler/brw/brw_shader.cpp +++ b/src/intel/compiler/brw/brw_shader.cpp @@ -17,6 +17,7 @@ #include "dev/intel_wa.h" #include "compiler/glsl_types.h" #include "compiler/nir/nir_builder.h" +#include "util/bitscan.h" #include "util/u_math.h" void @@ -999,7 +1000,17 @@ brw_cs_get_dispatch_info(const struct intel_device_info *devinfo, override_local_size ? override_local_size : prog_data->local_size; - const int simd = brw_simd_select_for_workgroup_size(devinfo, prog_data, sizes); + int simd = -1; + if (intel_use_jay(devinfo, MESA_SHADER_COMPUTE)) { + /* Currently Jay compiles only a single binary, just select that. In the + * future this needs to get smarter. + */ + assert(util_is_power_of_two_nonzero(prog_data->prog_mask)); + simd = util_logbase2(prog_data->prog_mask); + } else { + simd = brw_simd_select_for_workgroup_size(devinfo, prog_data, sizes); + } + assert(simd >= 0 && simd < 3); info.group_size = sizes[0] * sizes[1] * sizes[2];