From 4778fc1ab7b8a3fe44df82c7a0b48290a1dca1c2 Mon Sep 17 00:00:00 2001 From: Alyssa Rosenzweig Date: Thu, 26 Feb 2026 17:19:22 -0500 Subject: [PATCH] brw: disable hw generate local ID for jay Jay will need more work to handle these payloads properly especially in SIMD32. For now just disable the optimization for Jay for correctness. Signed-off-by: Alyssa Rosenzweig Reviewed-by: Lionel Landwerlin Reviewed-by: Kenneth Graunke Part-of: --- src/intel/compiler/brw/brw_nir_lower_cs_intrinsics.c | 4 +++- 1 file changed, 3 insertions(+), 1 deletion(-) diff --git a/src/intel/compiler/brw/brw_nir_lower_cs_intrinsics.c b/src/intel/compiler/brw/brw_nir_lower_cs_intrinsics.c index 7777a87ea84..befb01e0194 100644 --- a/src/intel/compiler/brw/brw_nir_lower_cs_intrinsics.c +++ b/src/intel/compiler/brw/brw_nir_lower_cs_intrinsics.c @@ -5,6 +5,7 @@ #include "brw_nir.h" #include "compiler/nir/nir_builder.h" +#include "dev/intel_debug.h" struct lower_intrinsics_state { nir_shader *nir; @@ -323,7 +324,8 @@ brw_nir_lower_cs_intrinsics(nir_shader *nir, nir->info.derivative_group != DERIVATIVE_GROUP_QUADS && !nir->info.workgroup_size_variable && util_is_power_of_two_nonzero(nir->info.workgroup_size[0]) && - util_is_power_of_two_nonzero(nir->info.workgroup_size[1])) { + util_is_power_of_two_nonzero(nir->info.workgroup_size[1]) && + !intel_use_jay(devinfo, nir->info.stage)) { state.hw_generated_local_id = true;