diff --git a/src/intel/compiler/brw_simd_selection.c b/src/intel/compiler/brw_simd_selection.c
index 2ecf568ad57..0c9bf1f4fd2 100644
--- a/src/intel/compiler/brw_simd_selection.c
+++ b/src/intel/compiler/brw_simd_selection.c
@@ -93,7 +93,13 @@ brw_simd_should_compile(void *mem_ctx,
          return false;
       }
 
-      /* TODO: Ignore SIMD larger than workgroup if previous SIMD already passed. */
+      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",
+            width, workgroup_size, width / 2);
+         return false;
+      }
 
       if (DIV_ROUND_UP(workgroup_size, width) > max_threads) {
          *error = ralloc_asprintf(
diff --git a/src/intel/compiler/test_simd_selection.cpp b/src/intel/compiler/test_simd_selection.cpp
index 7344c57aca4..6cbf9240a27 100644
--- a/src/intel/compiler/test_simd_selection.cpp
+++ b/src/intel/compiler/test_simd_selection.cpp
@@ -109,11 +109,10 @@ TEST_F(SIMDSelectionCS, WorkgroupSize1)
 
    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(SIMD16));
    ASSERT_FALSE(should_compile(SIMD32));
 
-   ASSERT_EQ(brw_simd_select(prog_data), SIMD16);
+   ASSERT_EQ(brw_simd_select(prog_data), SIMD8);
 }
 
 TEST_F(SIMDSelectionCS, WorkgroupSize8)
@@ -124,11 +123,10 @@ TEST_F(SIMDSelectionCS, WorkgroupSize8)
 
    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(SIMD16));
    ASSERT_FALSE(should_compile(SIMD32));
 
-   ASSERT_EQ(brw_simd_select(prog_data), SIMD16);
+   ASSERT_EQ(brw_simd_select(prog_data), SIMD8);
 }
 
 TEST_F(SIMDSelectionCS, WorkgroupSizeVariable)
@@ -147,7 +145,7 @@ TEST_F(SIMDSelectionCS, WorkgroupSizeVariable)
    ASSERT_EQ(prog_data->prog_mask, 1u << SIMD8 | 1u << SIMD16 | 1u << SIMD32);
 
    const unsigned wg_8_1_1[] = { 8, 1, 1 };
-   ASSERT_EQ(brw_simd_select_for_workgroup_size(devinfo, prog_data, wg_8_1_1), SIMD16);
+   ASSERT_EQ(brw_simd_select_for_workgroup_size(devinfo, prog_data, wg_8_1_1), SIMD8);
 
    const unsigned wg_16_1_1[] = { 16, 1, 1 };
    ASSERT_EQ(brw_simd_select_for_workgroup_size(devinfo, prog_data, wg_16_1_1), SIMD16);