intel/compiler: Create a struct to hold SIMD selection state

This is a preparation to decouple the storage of what SIMDs
compiled/spilled from the cs_prog_data.  This will allow reuse
of SIMD selection code by Bindless Shaders.

And since we have a struct now, move the error array there so
reduce the boilerplate of the users.

Reviewed-by: Lionel Landwerlin <lionel.g.landwerlin@intel.com>
Reviewed-by: Ivan Briano <ivan.briano@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/19601>

diff --git a/src/intel/compiler/brw_fs.cpp b/src/intel/compiler/brw_fs.cpp
index f52f8ca..c661a8c 100644 (file)
--- a/src/intel/compiler/brw_fs.cpp
+++ b/src/intel/compiler/brw_fs.cpp
@@ -7806,15 +7806,17 @@ brw_compile_cs(const struct brw_compiler *compiler,
       prog_data->local_size[2] = nir->info.workgroup_size[2];
    }
 
-   const unsigned required_dispatch_width =
-      brw_required_dispatch_width(&nir->info);
+   brw_simd_selection_state simd_state{
+      .mem_ctx = mem_ctx,
+      .devinfo = compiler->devinfo,
+      .prog_data = prog_data,
+      .required_width = brw_required_dispatch_width(&nir->info),
+   };
 
    std::unique_ptr<fs_visitor> v[3];
-   const char *error[3] = {0};
 
    for (unsigned simd = 0; simd < 3; simd++) {
-      if (!brw_simd_should_compile(mem_ctx, simd, compiler->devinfo, prog_data,
-                                   required_dispatch_width, &error[simd]))
+      if (!brw_simd_should_compile(simd_state, simd))
          continue;
 
       const unsigned dispatch_width = 8u << simd;
@@ -7847,9 +7849,9 @@ brw_compile_cs(const struct brw_compiler *compiler,
       if (v[simd]->run_cs(allow_spilling)) {
          cs_fill_push_const_info(compiler->devinfo, prog_data);
 
-         brw_simd_mark_compiled(simd, prog_data, v[simd]->spilled_any_registers);
+         brw_simd_mark_compiled(simd_state, simd, v[simd]->spilled_any_registers);
       } else {
-         error[simd] = ralloc_strdup(mem_ctx, v[simd]->fail_msg);
+         simd_state.error[simd] = ralloc_strdup(mem_ctx, v[simd]->fail_msg);
          if (simd > 0) {
             brw_shader_perf_log(compiler, params->log_data,
                                 "SIMD%u shader failed to compile: %s\n",
@@ -7858,10 +7860,11 @@ brw_compile_cs(const struct brw_compiler *compiler,
       }
    }
 
-   const int selected_simd = brw_simd_select(prog_data);
+   const int selected_simd = brw_simd_select(simd_state);
    if (selected_simd < 0) {
       params->error_str = ralloc_asprintf(mem_ctx, "Can't compile shader: %s, %s and %s.\n",
-                                          error[0], error[1], error[2]);;
+                                          simd_state.error[0], simd_state.error[1],
+                                          simd_state.error[2]);
       return NULL;
    }
 

diff --git a/src/intel/compiler/brw_mesh.cpp b/src/intel/compiler/brw_mesh.cpp
index f7c4b9e..64bcb52 100644 (file)
--- a/src/intel/compiler/brw_mesh.cpp
+++ b/src/intel/compiler/brw_mesh.cpp
@@ -265,15 +265,17 @@ brw_compile_task(const struct brw_compiler *compiler,
    prog_data->uses_drawid =
       BITSET_TEST(nir->info.system_values_read, SYSTEM_VALUE_DRAW_ID);
 
-   const unsigned required_dispatch_width =
-      brw_required_dispatch_width(&nir->info);
+   brw_simd_selection_state simd_state{
+      .mem_ctx = mem_ctx,
+      .devinfo = compiler->devinfo,
+      .prog_data = &prog_data->base,
+      .required_width = brw_required_dispatch_width(&nir->info),
+   };
 
    std::unique_ptr<fs_visitor> v[3];
-   const char *error[3] = {0};
 
    for (unsigned simd = 0; simd < 3; simd++) {
-      if (!brw_simd_should_compile(mem_ctx, simd, compiler->devinfo, &prog_data->base,
-                                   required_dispatch_width, &error[simd]))
+      if (!brw_simd_should_compile(simd_state, simd))
          continue;
 
       const unsigned dispatch_width = 8 << simd;
@@ -301,15 +303,16 @@ brw_compile_task(const struct brw_compiler *compiler,
       const bool allow_spilling = !prog_data->base.prog_mask;
 
       if (v[simd]->run_task(allow_spilling))
-         brw_simd_mark_compiled(simd, &prog_data->base, v[simd]->spilled_any_registers);
+         brw_simd_mark_compiled(simd_state, simd, v[simd]->spilled_any_registers);
       else
-         error[simd] = ralloc_strdup(mem_ctx, v[simd]->fail_msg);
+         simd_state.error[simd] = ralloc_strdup(mem_ctx, v[simd]->fail_msg);
    }
 
-   int selected_simd = brw_simd_select(&prog_data->base);
+   int selected_simd = brw_simd_select(simd_state);
    if (selected_simd < 0) {
       params->error_str = ralloc_asprintf(mem_ctx, "Can't compile shader: %s, %s and %s.\n",
-                                          error[0], error[1], error[2]);;
+                                          simd_state.error[0], simd_state.error[1],
+                                          simd_state.error[2]);
       return NULL;
    }
 
@@ -761,15 +764,17 @@ brw_compile_mesh(const struct brw_compiler *compiler,
    brw_compute_mue_map(nir, &prog_data->map);
    brw_nir_lower_mue_outputs(nir, &prog_data->map);
 
-   const unsigned required_dispatch_width =
-      brw_required_dispatch_width(&nir->info);
+   brw_simd_selection_state simd_state{
+      .mem_ctx = mem_ctx,
+      .devinfo = compiler->devinfo,
+      .prog_data = &prog_data->base,
+      .required_width = brw_required_dispatch_width(&nir->info),
+   };
 
    std::unique_ptr<fs_visitor> v[3];
-   const char *error[3] = {0};
 
    for (int simd = 0; simd < 3; simd++) {
-      if (!brw_simd_should_compile(mem_ctx, simd, compiler->devinfo, &prog_data->base,
-                                   required_dispatch_width, &error[simd]))
+      if (!brw_simd_should_compile(simd_state, simd))
          continue;
 
       const unsigned dispatch_width = 8 << simd;
@@ -809,15 +814,16 @@ brw_compile_mesh(const struct brw_compiler *compiler,
       const bool allow_spilling = !prog_data->base.prog_mask;
 
       if (v[simd]->run_mesh(allow_spilling))
-         brw_simd_mark_compiled(simd, &prog_data->base, v[simd]->spilled_any_registers);
+         brw_simd_mark_compiled(simd_state, simd, v[simd]->spilled_any_registers);
       else
-         error[simd] = ralloc_strdup(mem_ctx, v[simd]->fail_msg);
+         simd_state.error[simd] = ralloc_strdup(mem_ctx, v[simd]->fail_msg);
    }
 
-   int selected_simd = brw_simd_select(&prog_data->base);
+   int selected_simd = brw_simd_select(simd_state);
    if (selected_simd < 0) {
       params->error_str = ralloc_asprintf(mem_ctx, "Can't compile shader: %s, %s and %s.\n",
-                                          error[0], error[1], error[2]);;
+                                          simd_state.error[0], simd_state.error[1],
+                                          simd_state.error[2]);;
       return NULL;
    }
 

diff --git a/src/intel/compiler/brw_private.h b/src/intel/compiler/brw_private.h
index 70b6fd9..6f1374e 100644 (file)
--- a/src/intel/compiler/brw_private.h
+++ b/src/intel/compiler/brw_private.h
@@ -29,18 +29,22 @@
 
 unsigned brw_required_dispatch_width(const struct shader_info *info);
 
-bool brw_simd_should_compile(void *mem_ctx,
-                             unsigned simd,
-                             const struct intel_device_info *devinfo,
-                             struct brw_cs_prog_data *prog_data,
-                             unsigned required_dispatch_width,
-                             const char **error);
-
-void brw_simd_mark_compiled(unsigned simd,
-                            struct brw_cs_prog_data *prog_data,
-                            bool spilled);
-
-int brw_simd_select(const struct brw_cs_prog_data *prog_data);
+struct brw_simd_selection_state {
+   void *mem_ctx;
+   const struct intel_device_info *devinfo;
+
+   struct brw_cs_prog_data *prog_data;
+
+   unsigned required_width;
+
+   const char *error[3];
+};
+
+bool brw_simd_should_compile(brw_simd_selection_state &state, unsigned simd);
+
+void brw_simd_mark_compiled(brw_simd_selection_state &state, unsigned simd, bool spilled);
+
+int brw_simd_select(const brw_simd_selection_state &state);
 
 int brw_simd_select_for_workgroup_size(const struct intel_device_info *devinfo,
                                        const struct brw_cs_prog_data *prog_data,

diff --git a/src/intel/compiler/brw_simd_selection.cpp b/src/intel/compiler/brw_simd_selection.cpp
index 078ff0a..aa16674 100644 (file)
--- a/src/intel/compiler/brw_simd_selection.cpp
+++ b/src/intel/compiler/brw_simd_selection.cpp
@@ -47,16 +47,11 @@ test_bit(unsigned mask, unsigned bit) {
 }
 
 bool
-brw_simd_should_compile(void *mem_ctx,
-                        unsigned simd,
-                        const struct intel_device_info *devinfo,
-                        struct brw_cs_prog_data *prog_data,
-                        unsigned required,
-                        const char **error)
-
+brw_simd_should_compile(brw_simd_selection_state &state,
+                        unsigned simd)
 {
+   struct brw_cs_prog_data *prog_data = state.prog_data;
    assert(!test_bit(prog_data->prog_mask, simd));
-   assert(error);
 
    const unsigned width = 8u << simd;
 
@@ -68,8 +63,8 @@ brw_simd_should_compile(void *mem_ctx,
 
    if (!workgroup_size_variable) {
       if (test_bit(prog_data->prog_spilled, simd)) {
-         *error = ralloc_asprintf(
-            mem_ctx, "SIMD%u skipped because would spill", width);
+         state.error[simd] = ralloc_asprintf(
+            state.mem_ctx, "SIMD%u skipped because would spill", width);
          return false;
       }
 
@@ -77,26 +72,26 @@ brw_simd_should_compile(void *mem_ctx,
                                       prog_data->local_size[1] *
                                       prog_data->local_size[2];
 
-      unsigned max_threads = devinfo->max_cs_workgroup_threads;
+      unsigned max_threads = state.devinfo->max_cs_workgroup_threads;
 
-      if (required && required != width) {
-         *error = ralloc_asprintf(
-            mem_ctx, "SIMD%u skipped because required dispatch width is %u",
-            width, required);
+      if (state.required_width && state.required_width != width) {
+         state.error[simd] = ralloc_asprintf(
+            state.mem_ctx, "SIMD%u skipped because required dispatch width is %u",
+            width, state.required_width);
          return false;
       }
 
       if (simd > 0 && test_bit(prog_data->prog_mask, simd - 1) &&
           workgroup_size <= (width / 2)) {
-         *error = ralloc_asprintf(
-            mem_ctx, "SIMD%u skipped because workgroup size %u already fits in SIMD%u",
+         state.error[simd] = ralloc_asprintf(
+            state.mem_ctx, "SIMD%u skipped because workgroup size %u already fits in SIMD%u",
             width, workgroup_size, width / 2);
          return false;
       }
 
       if (DIV_ROUND_UP(workgroup_size, width) > max_threads) {
-         *error = ralloc_asprintf(
-            mem_ctx, "SIMD%u can't fit all %u invocations in %u threads",
+         state.error[simd] = ralloc_asprintf(
+            state.mem_ctx, "SIMD%u can't fit all %u invocations in %u threads",
             width, workgroup_size, max_threads);
          return false;
       }
@@ -107,23 +102,23 @@ brw_simd_should_compile(void *mem_ctx,
        */
       if (width == 32) {
          if (!INTEL_DEBUG(DEBUG_DO32) && prog_data->prog_mask) {
-            *error = ralloc_strdup(
-               mem_ctx, "SIMD32 skipped because not required");
+            state.error[simd] = ralloc_strdup(
+               state.mem_ctx, "SIMD32 skipped because not required");
             return false;
          }
       }
    }
 
    if (width == 32 && prog_data->base.ray_queries > 0) {
-      *error = ralloc_asprintf(
-         mem_ctx, "SIMD%u skipped because of ray queries",
+      state.error[simd] = ralloc_asprintf(
+         state.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",
+      state.error[simd] = ralloc_asprintf(
+         state.mem_ctx, "SIMD%u skipped because of bindless shader calls",
          width);
       return false;
    }
@@ -135,8 +130,8 @@ brw_simd_should_compile(void *mem_ctx,
    };
 
    if (unlikely(env_skip[simd])) {
-      *error = ralloc_asprintf(
-         mem_ctx, "SIMD%u skipped because INTEL_DEBUG=no%u",
+      state.error[simd] = ralloc_asprintf(
+         state.mem_ctx, "SIMD%u skipped because INTEL_DEBUG=no%u",
          width, width);
       return false;
    }
@@ -145,8 +140,9 @@ brw_simd_should_compile(void *mem_ctx,
 }
 
 void
-brw_simd_mark_compiled(unsigned simd, struct brw_cs_prog_data *prog_data, bool spilled)
+brw_simd_mark_compiled(brw_simd_selection_state &state, unsigned simd, bool spilled)
 {
+   struct brw_cs_prog_data *prog_data = state.prog_data;
    assert(!test_bit(prog_data->prog_mask, simd));
 
    prog_data->prog_mask |= 1u << simd;
@@ -159,8 +155,9 @@ brw_simd_mark_compiled(unsigned simd, struct brw_cs_prog_data *prog_data, bool s
 }
 
 int
-brw_simd_select(const struct brw_cs_prog_data *prog_data)
+brw_simd_select(const struct brw_simd_selection_state &state)
 {
+   const struct brw_cs_prog_data *prog_data = state.prog_data;
    assert((prog_data->prog_mask & ~0x7u) == 0);
    const unsigned not_spilled_mask =
       prog_data->prog_mask & ~prog_data->prog_spilled;
@@ -182,10 +179,12 @@ brw_simd_select_for_workgroup_size(const struct intel_device_info *devinfo,
 {
    if (!sizes || (prog_data->local_size[0] == sizes[0] &&
                   prog_data->local_size[1] == sizes[1] &&
-                  prog_data->local_size[2] == sizes[2]))
-      return brw_simd_select(prog_data);
-
-   void *mem_ctx = ralloc_context(NULL);
+                  prog_data->local_size[2] == sizes[2])) {
+      const brw_simd_selection_state simd_state{
+         .prog_data = const_cast<struct brw_cs_prog_data *>(prog_data),
+      };
+      return brw_simd_select(simd_state);
+   }
 
    struct brw_cs_prog_data cloned = *prog_data;
    for (unsigned i = 0; i < 3; i++)
@@ -194,20 +193,25 @@ brw_simd_select_for_workgroup_size(const struct intel_device_info *devinfo,
    cloned.prog_mask = 0;
    cloned.prog_spilled = 0;
 
-   const char *error[3] = {0};
+   void *mem_ctx = ralloc_context(NULL);
+
+   brw_simd_selection_state simd_state{
+      .mem_ctx = mem_ctx,
+      .devinfo = devinfo,
+      .prog_data = &cloned,
+   };
 
    for (unsigned simd = 0; simd < 3; simd++) {
       /* We are not recompiling, so use original results of prog_mask and
        * prog_spilled as they will already contain all possible compilations.
        */
-      if (brw_simd_should_compile(mem_ctx, simd, devinfo, &cloned,
-                                  0 /* required_dispatch_width */, &error[simd]) &&
+      if (brw_simd_should_compile(simd_state, simd) &&
           test_bit(prog_data->prog_mask, simd)) {
-         brw_simd_mark_compiled(simd, &cloned, test_bit(prog_data->prog_spilled, simd));
+         brw_simd_mark_compiled(simd_state, simd, test_bit(prog_data->prog_spilled, simd));
       }
    }
 
    ralloc_free(mem_ctx);
 
-   return brw_simd_select(&cloned);
+   return brw_simd_select(simd_state);
 }

diff --git a/src/intel/compiler/test_simd_selection.cpp b/src/intel/compiler/test_simd_selection.cpp
index 5783c32..c8a6f02 100644 (file)
--- a/src/intel/compiler/test_simd_selection.cpp
+++ b/src/intel/compiler/test_simd_selection.cpp
@@ -41,27 +41,26 @@ const bool not_spilled = false;
 
 class SIMDSelectionTest : public ::testing::Test {
 protected:
-   SIMDSelectionTest() : error{NULL, NULL, NULL} {
-      mem_ctx = ralloc_context(NULL);
-      devinfo = rzalloc(mem_ctx, intel_device_info);
-      prog_data = rzalloc(mem_ctx, struct brw_cs_prog_data);
-      required_dispatch_width = 0;
+   SIMDSelectionTest()
+   : mem_ctx(ralloc_context(NULL))
+   , devinfo(rzalloc(mem_ctx, intel_device_info))
+   , prog_data(rzalloc(mem_ctx, struct brw_cs_prog_data))
+   , simd_state{
+      .mem_ctx = mem_ctx,
+      .devinfo = devinfo,
+      .prog_data = prog_data,
+     }
+   {
    }
 
    ~SIMDSelectionTest() {
       ralloc_free(mem_ctx);
    };
 
-   bool should_compile(unsigned simd) {
-      return brw_simd_should_compile(mem_ctx, simd, devinfo, prog_data,
-                                     required_dispatch_width, &error[simd]);
-   }
-
    void *mem_ctx;
    intel_device_info *devinfo;
    struct brw_cs_prog_data *prog_data;
-   const char *error[3];
-   unsigned required_dispatch_width;
+   brw_simd_selection_state simd_state;
 };
 
 class SIMDSelectionCS : public SIMDSelectionTest {
@@ -78,13 +77,13 @@ protected:
 
 TEST_F(SIMDSelectionCS, DefaultsToSIMD16)
 {
-   ASSERT_TRUE(should_compile(SIMD8));
-   brw_simd_mark_compiled(SIMD8, prog_data, not_spilled);
-   ASSERT_TRUE(should_compile(SIMD16));
-   brw_simd_mark_compiled(SIMD16, prog_data, not_spilled);
-   ASSERT_FALSE(should_compile(SIMD32));
+   ASSERT_TRUE(brw_simd_should_compile(simd_state, SIMD8));
+   brw_simd_mark_compiled(simd_state, SIMD8, not_spilled);
+   ASSERT_TRUE(brw_simd_should_compile(simd_state, SIMD16));
+   brw_simd_mark_compiled(simd_state, SIMD16, not_spilled);
+   ASSERT_FALSE(brw_simd_should_compile(simd_state, SIMD32));
 
-   ASSERT_EQ(brw_simd_select(prog_data), SIMD16);
+   ASSERT_EQ(brw_simd_select(simd_state), SIMD16);
 }
 
 TEST_F(SIMDSelectionCS, TooBigFor16)
@@ -93,12 +92,12 @@ TEST_F(SIMDSelectionCS, TooBigFor16)
    prog_data->local_size[1] = 32;
    prog_data->local_size[2] = 1;
 
-   ASSERT_FALSE(should_compile(SIMD8));
-   ASSERT_FALSE(should_compile(SIMD16));
-   ASSERT_TRUE(should_compile(SIMD32));
-   brw_simd_mark_compiled(SIMD32, prog_data, spilled);
+   ASSERT_FALSE(brw_simd_should_compile(simd_state, SIMD8));
+   ASSERT_FALSE(brw_simd_should_compile(simd_state, SIMD16));
+   ASSERT_TRUE(brw_simd_should_compile(simd_state, SIMD32));
+   brw_simd_mark_compiled(simd_state, SIMD32, spilled);
 
-   ASSERT_EQ(brw_simd_select(prog_data), SIMD32);
+   ASSERT_EQ(brw_simd_select(simd_state), SIMD32);
 }
 
 TEST_F(SIMDSelectionCS, WorkgroupSize1)
@@ -107,12 +106,12 @@ TEST_F(SIMDSelectionCS, WorkgroupSize1)
    prog_data->local_size[1] = 1;
    prog_data->local_size[2] = 1;
 
-   ASSERT_TRUE(should_compile(SIMD8));
-   brw_simd_mark_compiled(SIMD8, prog_data, not_spilled);
-   ASSERT_FALSE(should_compile(SIMD16));
-   ASSERT_FALSE(should_compile(SIMD32));
+   ASSERT_TRUE(brw_simd_should_compile(simd_state, SIMD8));
+   brw_simd_mark_compiled(simd_state, SIMD8, not_spilled);
+   ASSERT_FALSE(brw_simd_should_compile(simd_state, SIMD16));
+   ASSERT_FALSE(brw_simd_should_compile(simd_state, SIMD32));
 
-   ASSERT_EQ(brw_simd_select(prog_data), SIMD8);
+   ASSERT_EQ(brw_simd_select(simd_state), SIMD8);
 }
 
 TEST_F(SIMDSelectionCS, WorkgroupSize8)
@@ -121,12 +120,12 @@ TEST_F(SIMDSelectionCS, WorkgroupSize8)
    prog_data->local_size[1] = 1;
    prog_data->local_size[2] = 1;
 
-   ASSERT_TRUE(should_compile(SIMD8));
-   brw_simd_mark_compiled(SIMD8, prog_data, not_spilled);
-   ASSERT_FALSE(should_compile(SIMD16));
-   ASSERT_FALSE(should_compile(SIMD32));
+   ASSERT_TRUE(brw_simd_should_compile(simd_state, SIMD8));
+   brw_simd_mark_compiled(simd_state, SIMD8, not_spilled);
+   ASSERT_FALSE(brw_simd_should_compile(simd_state, SIMD16));
+   ASSERT_FALSE(brw_simd_should_compile(simd_state, SIMD32));
 
-   ASSERT_EQ(brw_simd_select(prog_data), SIMD8);
+   ASSERT_EQ(brw_simd_select(simd_state), SIMD8);
 }
 
 TEST_F(SIMDSelectionCS, WorkgroupSizeVariable)
@@ -135,12 +134,12 @@ TEST_F(SIMDSelectionCS, WorkgroupSizeVariable)
    prog_data->local_size[1] = 0;
    prog_data->local_size[2] = 0;
 
-   ASSERT_TRUE(should_compile(SIMD8));
-   brw_simd_mark_compiled(SIMD8, prog_data, not_spilled);
-   ASSERT_TRUE(should_compile(SIMD16));
-   brw_simd_mark_compiled(SIMD16, prog_data, not_spilled);
-   ASSERT_TRUE(should_compile(SIMD32));
-   brw_simd_mark_compiled(SIMD32, prog_data, not_spilled);
+   ASSERT_TRUE(brw_simd_should_compile(simd_state, SIMD8));
+   brw_simd_mark_compiled(simd_state, SIMD8, not_spilled);
+   ASSERT_TRUE(brw_simd_should_compile(simd_state, SIMD16));
+   brw_simd_mark_compiled(simd_state, SIMD16, not_spilled);
+   ASSERT_TRUE(brw_simd_should_compile(simd_state, SIMD32));
+   brw_simd_mark_compiled(simd_state, SIMD32, not_spilled);
 
    ASSERT_EQ(prog_data->prog_mask, 1u << SIMD8 | 1u << SIMD16 | 1u << SIMD32);
 
@@ -160,12 +159,12 @@ TEST_F(SIMDSelectionCS, WorkgroupSizeVariableSpilled)
    prog_data->local_size[1] = 0;
    prog_data->local_size[2] = 0;
 
-   ASSERT_TRUE(should_compile(SIMD8));
-   brw_simd_mark_compiled(SIMD8, prog_data, spilled);
-   ASSERT_TRUE(should_compile(SIMD16));
-   brw_simd_mark_compiled(SIMD16, prog_data, spilled);
-   ASSERT_TRUE(should_compile(SIMD32));
-   brw_simd_mark_compiled(SIMD32, prog_data, spilled);
+   ASSERT_TRUE(brw_simd_should_compile(simd_state, SIMD8));
+   brw_simd_mark_compiled(simd_state, SIMD8, spilled);
+   ASSERT_TRUE(brw_simd_should_compile(simd_state, SIMD16));
+   brw_simd_mark_compiled(simd_state, SIMD16, spilled);
+   ASSERT_TRUE(brw_simd_should_compile(simd_state, SIMD32));
+   brw_simd_mark_compiled(simd_state, SIMD32, spilled);
 
    ASSERT_EQ(prog_data->prog_mask, 1u << SIMD8 | 1u << SIMD16 | 1u << SIMD32);
 
@@ -185,11 +184,11 @@ TEST_F(SIMDSelectionCS, WorkgroupSizeVariableNoSIMD8)
    prog_data->local_size[1] = 0;
    prog_data->local_size[2] = 0;
 
-   ASSERT_TRUE(should_compile(SIMD8));
-   ASSERT_TRUE(should_compile(SIMD16));
-   brw_simd_mark_compiled(SIMD16, prog_data, not_spilled);
-   ASSERT_TRUE(should_compile(SIMD32));
-   brw_simd_mark_compiled(SIMD32, prog_data, not_spilled);
+   ASSERT_TRUE(brw_simd_should_compile(simd_state, SIMD8));
+   ASSERT_TRUE(brw_simd_should_compile(simd_state, SIMD16));
+   brw_simd_mark_compiled(simd_state, SIMD16, not_spilled);
+   ASSERT_TRUE(brw_simd_should_compile(simd_state, SIMD32));
+   brw_simd_mark_compiled(simd_state, SIMD32, not_spilled);
 
    ASSERT_EQ(prog_data->prog_mask, 1u << SIMD16 | 1u << SIMD32);
 
@@ -209,11 +208,11 @@ TEST_F(SIMDSelectionCS, WorkgroupSizeVariableNoSIMD16)
    prog_data->local_size[1] = 0;
    prog_data->local_size[2] = 0;
 
-   ASSERT_TRUE(should_compile(SIMD8));
-   brw_simd_mark_compiled(SIMD8, prog_data, not_spilled);
-   ASSERT_TRUE(should_compile(SIMD16));
-   ASSERT_TRUE(should_compile(SIMD32));
-   brw_simd_mark_compiled(SIMD32, prog_data, not_spilled);
+   ASSERT_TRUE(brw_simd_should_compile(simd_state, SIMD8));
+   brw_simd_mark_compiled(simd_state, SIMD8, not_spilled);
+   ASSERT_TRUE(brw_simd_should_compile(simd_state, SIMD16));
+   ASSERT_TRUE(brw_simd_should_compile(simd_state, SIMD32));
+   brw_simd_mark_compiled(simd_state, SIMD32, not_spilled);
 
    ASSERT_EQ(prog_data->prog_mask, 1u << SIMD8 | 1u << SIMD32);
 
@@ -233,10 +232,10 @@ TEST_F(SIMDSelectionCS, WorkgroupSizeVariableNoSIMD8NoSIMD16)
    prog_data->local_size[1] = 0;
    prog_data->local_size[2] = 0;
 
-   ASSERT_TRUE(should_compile(SIMD8));
-   ASSERT_TRUE(should_compile(SIMD16));
-   ASSERT_TRUE(should_compile(SIMD32));
-   brw_simd_mark_compiled(SIMD32, prog_data, not_spilled);
+   ASSERT_TRUE(brw_simd_should_compile(simd_state, SIMD8));
+   ASSERT_TRUE(brw_simd_should_compile(simd_state, SIMD16));
+   ASSERT_TRUE(brw_simd_should_compile(simd_state, SIMD32));
+   brw_simd_mark_compiled(simd_state, SIMD32, not_spilled);
 
    ASSERT_EQ(prog_data->prog_mask, 1u << SIMD32);
 
@@ -252,118 +251,118 @@ TEST_F(SIMDSelectionCS, WorkgroupSizeVariableNoSIMD8NoSIMD16)
 
 TEST_F(SIMDSelectionCS, SpillAtSIMD8)
 {
-   ASSERT_TRUE(should_compile(SIMD8));
-   brw_simd_mark_compiled(SIMD8, prog_data, spilled);
-   ASSERT_FALSE(should_compile(SIMD16));
-   ASSERT_FALSE(should_compile(SIMD32));
+   ASSERT_TRUE(brw_simd_should_compile(simd_state, SIMD8));
+   brw_simd_mark_compiled(simd_state, SIMD8, spilled);
+   ASSERT_FALSE(brw_simd_should_compile(simd_state, SIMD16));
+   ASSERT_FALSE(brw_simd_should_compile(simd_state, SIMD32));
 
-   ASSERT_EQ(brw_simd_select(prog_data), SIMD8);
+   ASSERT_EQ(brw_simd_select(simd_state), SIMD8);
 }
 
 TEST_F(SIMDSelectionCS, SpillAtSIMD16)
 {
-   ASSERT_TRUE(should_compile(SIMD8));
-   brw_simd_mark_compiled(SIMD8, prog_data, not_spilled);
-   ASSERT_TRUE(should_compile(SIMD16));
-   brw_simd_mark_compiled(SIMD16, prog_data, spilled);
-   ASSERT_FALSE(should_compile(SIMD32));
+   ASSERT_TRUE(brw_simd_should_compile(simd_state, SIMD8));
+   brw_simd_mark_compiled(simd_state, SIMD8, not_spilled);
+   ASSERT_TRUE(brw_simd_should_compile(simd_state, SIMD16));
+   brw_simd_mark_compiled(simd_state, SIMD16, spilled);
+   ASSERT_FALSE(brw_simd_should_compile(simd_state, SIMD32));
 
-   ASSERT_EQ(brw_simd_select(prog_data), SIMD8);
+   ASSERT_EQ(brw_simd_select(simd_state), SIMD8);
 }
 
 TEST_F(SIMDSelectionCS, EnvironmentVariable32)
 {
    intel_debug |= DEBUG_DO32;
 
-   ASSERT_TRUE(should_compile(SIMD8));
-   brw_simd_mark_compiled(SIMD8, prog_data, not_spilled);
-   ASSERT_TRUE(should_compile(SIMD16));
-   brw_simd_mark_compiled(SIMD16, prog_data, not_spilled);
-   ASSERT_TRUE(should_compile(SIMD32));
-   brw_simd_mark_compiled(SIMD32, prog_data, not_spilled);
+   ASSERT_TRUE(brw_simd_should_compile(simd_state, SIMD8));
+   brw_simd_mark_compiled(simd_state, SIMD8, not_spilled);
+   ASSERT_TRUE(brw_simd_should_compile(simd_state, SIMD16));
+   brw_simd_mark_compiled(simd_state, SIMD16, not_spilled);
+   ASSERT_TRUE(brw_simd_should_compile(simd_state, SIMD32));
+   brw_simd_mark_compiled(simd_state, SIMD32, not_spilled);
 
-   ASSERT_EQ(brw_simd_select(prog_data), SIMD32);
+   ASSERT_EQ(brw_simd_select(simd_state), SIMD32);
 }
 
 TEST_F(SIMDSelectionCS, EnvironmentVariable32ButSpills)
 {
    intel_debug |= DEBUG_DO32;
 
-   ASSERT_TRUE(should_compile(SIMD8));
-   brw_simd_mark_compiled(SIMD8, prog_data, not_spilled);
-   ASSERT_TRUE(should_compile(SIMD16));
-   brw_simd_mark_compiled(SIMD16, prog_data, not_spilled);
-   ASSERT_TRUE(should_compile(SIMD32));
-   brw_simd_mark_compiled(SIMD32, prog_data, spilled);
+   ASSERT_TRUE(brw_simd_should_compile(simd_state, SIMD8));
+   brw_simd_mark_compiled(simd_state, SIMD8, not_spilled);
+   ASSERT_TRUE(brw_simd_should_compile(simd_state, SIMD16));
+   brw_simd_mark_compiled(simd_state, SIMD16, not_spilled);
+   ASSERT_TRUE(brw_simd_should_compile(simd_state, SIMD32));
+   brw_simd_mark_compiled(simd_state, SIMD32, spilled);
 
-   ASSERT_EQ(brw_simd_select(prog_data), SIMD16);
+   ASSERT_EQ(brw_simd_select(simd_state), SIMD16);
 }
 
 TEST_F(SIMDSelectionCS, Require8)
 {
-   required_dispatch_width = 8;
+   simd_state.required_width = 8;
 
-   ASSERT_TRUE(should_compile(SIMD8));
-   brw_simd_mark_compiled(SIMD8, prog_data, not_spilled);
-   ASSERT_FALSE(should_compile(SIMD16));
-   ASSERT_FALSE(should_compile(SIMD32));
+   ASSERT_TRUE(brw_simd_should_compile(simd_state, SIMD8));
+   brw_simd_mark_compiled(simd_state, SIMD8, not_spilled);
+   ASSERT_FALSE(brw_simd_should_compile(simd_state, SIMD16));
+   ASSERT_FALSE(brw_simd_should_compile(simd_state, SIMD32));
 
-   ASSERT_EQ(brw_simd_select(prog_data), SIMD8);
+   ASSERT_EQ(brw_simd_select(simd_state), SIMD8);
 }
 
 TEST_F(SIMDSelectionCS, Require8ErrorWhenNotCompile)
 {
-   required_dispatch_width = 8;
+   simd_state.required_width = 8;
 
-   ASSERT_TRUE(should_compile(SIMD8));
-   ASSERT_FALSE(should_compile(SIMD16));
-   ASSERT_FALSE(should_compile(SIMD32));
+   ASSERT_TRUE(brw_simd_should_compile(simd_state, SIMD8));
+   ASSERT_FALSE(brw_simd_should_compile(simd_state, SIMD16));
+   ASSERT_FALSE(brw_simd_should_compile(simd_state, SIMD32));
 
-   ASSERT_EQ(brw_simd_select(prog_data), -1);
+   ASSERT_EQ(brw_simd_select(simd_state), -1);
 }
 
 TEST_F(SIMDSelectionCS, Require16)
 {
-   required_dispatch_width = 16;
+   simd_state.required_width = 16;
 
-   ASSERT_FALSE(should_compile(SIMD8));
-   ASSERT_TRUE(should_compile(SIMD16));
-   brw_simd_mark_compiled(SIMD16, prog_data, not_spilled);
-   ASSERT_FALSE(should_compile(SIMD32));
+   ASSERT_FALSE(brw_simd_should_compile(simd_state, SIMD8));
+   ASSERT_TRUE(brw_simd_should_compile(simd_state, SIMD16));
+   brw_simd_mark_compiled(simd_state, SIMD16, not_spilled);
+   ASSERT_FALSE(brw_simd_should_compile(simd_state, SIMD32));
 
-   ASSERT_EQ(brw_simd_select(prog_data), SIMD16);
+   ASSERT_EQ(brw_simd_select(simd_state), SIMD16);
 }
 
 TEST_F(SIMDSelectionCS, Require16ErrorWhenNotCompile)
 {
-   required_dispatch_width = 16;
+   simd_state.required_width = 16;
 
-   ASSERT_FALSE(should_compile(SIMD8));
-   ASSERT_TRUE(should_compile(SIMD16));
-   ASSERT_FALSE(should_compile(SIMD32));
+   ASSERT_FALSE(brw_simd_should_compile(simd_state, SIMD8));
+   ASSERT_TRUE(brw_simd_should_compile(simd_state, SIMD16));
+   ASSERT_FALSE(brw_simd_should_compile(simd_state, SIMD32));
 
-   ASSERT_EQ(brw_simd_select(prog_data), -1);
+   ASSERT_EQ(brw_simd_select(simd_state), -1);
 }
 
 TEST_F(SIMDSelectionCS, Require32)
 {
-   required_dispatch_width = 32;
+   simd_state.required_width = 32;
 
-   ASSERT_FALSE(should_compile(SIMD8));
-   ASSERT_FALSE(should_compile(SIMD16));
-   ASSERT_TRUE(should_compile(SIMD32));
-   brw_simd_mark_compiled(SIMD32, prog_data, not_spilled);
+   ASSERT_FALSE(brw_simd_should_compile(simd_state, SIMD8));
+   ASSERT_FALSE(brw_simd_should_compile(simd_state, SIMD16));
+   ASSERT_TRUE(brw_simd_should_compile(simd_state, SIMD32));
+   brw_simd_mark_compiled(simd_state, SIMD32, not_spilled);
 
-   ASSERT_EQ(brw_simd_select(prog_data), SIMD32);
+   ASSERT_EQ(brw_simd_select(simd_state), SIMD32);
 }
 
 TEST_F(SIMDSelectionCS, Require32ErrorWhenNotCompile)
 {
-   required_dispatch_width = 32;
+   simd_state.required_width = 32;
 
-   ASSERT_FALSE(should_compile(SIMD8));
-   ASSERT_FALSE(should_compile(SIMD16));
-   ASSERT_TRUE(should_compile(SIMD32));
+   ASSERT_FALSE(brw_simd_should_compile(simd_state, SIMD8));
+   ASSERT_FALSE(brw_simd_should_compile(simd_state, SIMD16));
+   ASSERT_TRUE(brw_simd_should_compile(simd_state, SIMD32));
 
-   ASSERT_EQ(brw_simd_select(prog_data), -1);
+   ASSERT_EQ(brw_simd_select(simd_state), -1);
 }