From 7397ba61c2d3e6f0347c34e863ca2ebd971f098d Mon Sep 17 00:00:00 2001
From: Francisco Jerez <currojerez@riseup.net>
Date: Fri, 1 Dec 2023 20:35:31 -0800
Subject: [PATCH] intel/fs/xe2+: Stop building SIMD8 compute-like shaders
 (CS/BS/TS/MS).

SIMD8 kernels are no longer able to utilize the ALUs efficiently,
since they have twice the vector width as previous platforms.  However
even though there aren't many reasons to use it, SIMD8 is still
supported by the instruction set technically, and it will still be
used for some SIMD-lowering sequences.

Reviewed-by: Caio Oliveira <caio.oliveira@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/26605>
---
 src/intel/compiler/brw_fs.cpp             |  3 +++
 src/intel/compiler/brw_simd_selection.cpp | 10 ++++++++--
 2 files changed, 11 insertions(+), 2 deletions(-)

diff --git a/src/intel/compiler/brw_fs.cpp b/src/intel/compiler/brw_fs.cpp
index 94aa4b8013b..6f28f54c7d0 100644
--- a/src/intel/compiler/brw_fs.cpp
+++ b/src/intel/compiler/brw_fs.cpp
@@ -8126,6 +8126,9 @@ compile_single_bs(const struct brw_compiler *compiler,
 
       const unsigned dispatch_width = 8u << simd;
 
+      if (dispatch_width == 8 && compiler->devinfo->ver >= 20)
+         continue;
+
       v[simd] = std::make_unique<fs_visitor>(compiler, &params->base,
                                              &key->base,
                                              &prog_data->base, shader,
diff --git a/src/intel/compiler/brw_simd_selection.cpp b/src/intel/compiler/brw_simd_selection.cpp
index 09d1646da72..05f9394c0d4 100644
--- a/src/intel/compiler/brw_simd_selection.cpp
+++ b/src/intel/compiler/brw_simd_selection.cpp
@@ -104,7 +104,8 @@ brw_simd_should_compile(brw_simd_selection_state &state, unsigned simd)
 
          unsigned max_threads = state.devinfo->max_cs_workgroup_threads;
 
-         if (simd > 0 && state.compiled[simd - 1] &&
+         const unsigned min_simd = state.devinfo->ver >= 20 ? 1 : 0;
+         if (simd > min_simd && state.compiled[simd - 1] &&
             workgroup_size <= (width / 2)) {
             state.error[simd] = "Workgroup size already fits in smaller SIMD";
             return false;
@@ -120,7 +121,7 @@ brw_simd_should_compile(brw_simd_selection_state &state, unsigned simd)
        *
        * TODO: Use performance_analysis and drop this rule.
        */
-      if (width == 32) {
+      if (width == 32 && state.devinfo->ver < 20) {
          if (!INTEL_DEBUG(DEBUG_DO32) && (state.compiled[0] || state.compiled[1])) {
             state.error[simd] = "SIMD32 not required (use INTEL_DEBUG=do32 to force)";
             return false;
@@ -128,6 +129,11 @@ brw_simd_should_compile(brw_simd_selection_state &state, unsigned simd)
       }
    }
 
+   if (width == 8 && state.devinfo->ver >= 20) {
+      state.error[simd] = "SIMD8 not supported on Xe2+";
+      return false;
+   }
+
    if (width == 32 && cs_prog_data && cs_prog_data->base.ray_queries > 0) {
       state.error[simd] = "Ray queries not supported";
       return false;