diff --git a/test/smoke-dev/big-jump-loop-split-2/big_jump_loop_split_2.c b/test/smoke-dev/big-jump-loop-split-2/big_jump_loop_split_2.c
index 975d38af4e..e8a17ff7fa 100644
--- a/test/smoke-dev/big-jump-loop-split-2/big_jump_loop_split_2.c
+++ b/test/smoke-dev/big-jump-loop-split-2/big_jump_loop_split_2.c
@@ -105,13 +105,15 @@ int main()
   return rc;
 }
 
-/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:7 ConstWGSize:256  args: 6 teamsXthrds:([[S:[ ]*]][[NUM_TEAMS:[0-9]+]]X  32)
-/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:7 ConstWGSize:64  args: 6 teamsXthrds:([[S:[ ]*]][[NUM_TEAMS:[0-9]+]]X  64)
-/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:7 ConstWGSize:256  args: 6 teamsXthrds:([[S:[ ]*]][[NUM_TEAMS:[0-9]+]]X  32)
-/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:7 ConstWGSize:128  args: 6 teamsXthrds:([[S:[ ]*]][[NUM_TEAMS:[0-9]+]]X  32)
-/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:7 ConstWGSize:256  args: 6 teamsXthrds:([[S:[ ]*]][[NUM_TEAMS:[0-9]+]]X  32)
+// Every loop here is small. Its kernel gets the full block size (ConstWGSize,
+// or the thread_limit clause) and as many teams as its trip count needs.
+/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:7 ConstWGSize:256  args: 6 teamsXthrds:({{ *}}1X 256)
+/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:7 ConstWGSize:64  args: 6 teamsXthrds:({{ *}}1X  64)
+/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:7 ConstWGSize:256  args: 6 teamsXthrds:({{ *}}1X 256)
+/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:7 ConstWGSize:128  args: 6 teamsXthrds:({{ *}}1X 128)
+/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:7 ConstWGSize:256  args: 6 teamsXthrds:({{ *}}1X 256)
 /// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:3 ConstWGSize:{{288|320}}  args: 8 teamsXthrds:([[S:[ ]*]][[NUM_TEAMS:[0-9]+]]X 256)
-/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:7 ConstWGSize:128  args: 6 teamsXthrds:([[S:[ ]*]][[NUM_TEAMS:[0-9]+]]X 128)
-/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:7 ConstWGSize:256  args: 6 teamsXthrds:([[S:[ ]*]][[NUM_TEAMS:[0-9]+]]X  32)
-/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:7 ConstWGSize:256  args: 6 teamsXthrds:([[S:[ ]*]][[NUM_TEAMS:[0-9]+]]X  32)
+/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:7 ConstWGSize:128  args: 6 teamsXthrds:({{ *}}1X 128)
+/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:7 ConstWGSize:256  args: 6 teamsXthrds:({{ *}}1X 256)
+/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:7 ConstWGSize:256  args: 6 teamsXthrds:({{ *}}1X 256)
 
diff --git a/test/smoke-dev/big-jump-loop-split-4/big_jump_loop_split_4.c b/test/smoke-dev/big-jump-loop-split-4/big_jump_loop_split_4.c
index 0e87d6e19c..e15430bc0e 100644
--- a/test/smoke-dev/big-jump-loop-split-4/big_jump_loop_split_4.c
+++ b/test/smoke-dev/big-jump-loop-split-4/big_jump_loop_split_4.c
@@ -129,17 +129,21 @@ int main()
   return rc;
 }
 
-/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:7 ConstWGSize:1024  args: 6 teamsXthrds:([[S:[ ]*]][[NUM_TEAMS:[0-9]+]]X 128)
-/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:7 ConstWGSize:1024  args: 6 teamsXthrds:([[S:[ ]*]][[NUM_TEAMS:[0-9]+]]X1024)
-/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:7 ConstWGSize:1024  args: 6 teamsXthrds:([[S:[ ]*]][[NUM_TEAMS:[0-9]+]]X 768)
-/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:7 ConstWGSize:1024  args: 6 teamsXthrds:([[S:[ ]*]][[NUM_TEAMS:[0-9]+]]X1024)
-/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:7 ConstWGSize:1024  args: 6 teamsXthrds:([[S:[ ]*]][[NUM_TEAMS:[0-9]+]]X1024)
-/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:7 ConstWGSize:1024  args: 6 teamsXthrds:([[S:[ ]*]][[NUM_TEAMS:[0-9]+]]X1024)
-/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:3 ConstWGSize:{{288|320}}  args: 8 teamsXthrds:([[S:[ ]*]][[NUM_TEAMS:[0-9]+]]X 128)
-/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:7 ConstWGSize:1024  args: 6 teamsXthrds:([[S:[ ]*]][[NUM_TEAMS:[0-9]+]]X 512)
-/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:7 ConstWGSize:1024  args: 6 teamsXthrds:([[S:[ ]*]][[NUM_TEAMS:[0-9]+]]X1024)
-/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:7 ConstWGSize:1024  args: 6 teamsXthrds:([[S:[ ]*]][[NUM_TEAMS:[0-9]+]]X  64)
-/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:7 ConstWGSize:1024  args: 6 teamsXthrds:([[S:[ ]*]][[NUM_TEAMS:[0-9]+]]X1024)
-/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:7 ConstWGSize:1024  args: 6 teamsXthrds:([[S:[ ]*]][[NUM_TEAMS:[0-9]+]]X1024)
-/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:7 ConstWGSize:1024  args: 6 teamsXthrds:([[S:[ ]*]][[NUM_TEAMS:[0-9]+]]X  32)
+// Every kernel gets the full block size (ConstWGSize, or the thread_limit or
+// num_threads clause) and as many teams as its trip count needs, at most
+// num_teams. Where num_teams asks for more teams than needed, the runtime may
+// launch either.
+/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:7 ConstWGSize:1024  args: 6 teamsXthrds:({{ *}}5X 128)
+/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:7 ConstWGSize:1024  args: 6 teamsXthrds:({{ *}}{{1[0-9]|20}}X1024)
+/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:7 ConstWGSize:1024  args: 6 teamsXthrds:({{ *}}{{1[0-9]|20}}X 768)
+/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:7 ConstWGSize:1024  args: 6 teamsXthrds:({{ *}}10X1024)
+/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:7 ConstWGSize:1024  args: 6 teamsXthrds:({{ *}}10X1024)
+/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:7 ConstWGSize:1024  args: 6 teamsXthrds:({{ *}}10X1024)
+/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:3 ConstWGSize:{{288|320}}  args: 8 teamsXthrds:({{ *}}20X 128)
+/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:7 ConstWGSize:1024  args: 6 teamsXthrds:({{ *}}20X 512)
+/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:7 ConstWGSize:1024  args: 6 teamsXthrds:({{ *}}7X1024)
+/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:7 ConstWGSize:1024  args: 6 teamsXthrds:({{ *}}20X  64)
+/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:7 ConstWGSize:1024  args: 6 teamsXthrds:({{ *}}10X1024)
+/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:7 ConstWGSize:1024  args: 6 teamsXthrds:({{ *}}10X1024)
+/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:7 ConstWGSize:1024  args: 6 teamsXthrds:({{ *}}3X1024)
 
diff --git a/test/smoke-dev/no-loop-split-2/no_loop_split_2.c b/test/smoke-dev/no-loop-split-2/no_loop_split_2.c
index 45c43106f9..9e0ce14f64 100644
--- a/test/smoke-dev/no-loop-split-2/no_loop_split_2.c
+++ b/test/smoke-dev/no-loop-split-2/no_loop_split_2.c
@@ -97,13 +97,15 @@ int main()
   return rc;
 }
 
-/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:6 ConstWGSize:64  args: 6 teamsXthrds:([[S:[ ]*]][[NUM_TEAMS:[0-9]+]]X 64)
-/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:6 ConstWGSize:100  args: 6 teamsXthrds:([[S:[ ]*]][[NUM_TEAMS:[0-9]+]]X 100)
-/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:6 ConstWGSize:256  args: 6 teamsXthrds:([[S:[ ]*]][[NUM_TEAMS:[0-9]+]]X  32)
-/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:6 ConstWGSize:256  args: 6 teamsXthrds:([[S:[ ]*]][[NUM_TEAMS:[0-9]+]]X  32)
-/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:6 ConstWGSize:256  args: 6 teamsXthrds:([[S:[ ]*]][[NUM_TEAMS:[0-9]+]]X  32)
+// Every loop here is small. Its kernel gets the full block size (ConstWGSize,
+// or the thread_limit clause) and as many teams as its trip count needs.
+/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:6 ConstWGSize:64  args: 6 teamsXthrds:({{ *}}1X  64)
+/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:6 ConstWGSize:100  args: 6 teamsXthrds:({{ *}}1X 100)
+/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:6 ConstWGSize:256  args: 6 teamsXthrds:({{ *}}1X 256)
+/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:6 ConstWGSize:256  args: 6 teamsXthrds:({{ *}}1X 256)
+/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:6 ConstWGSize:256  args: 6 teamsXthrds:({{ *}}1X 256)
 /// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:3 ConstWGSize:{{288|320}}  args: 8 teamsXthrds:([[S:[ ]*]][[NUM_TEAMS:[0-9]+]]X 256)
-/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:6 ConstWGSize:128  args: 6 teamsXthrds:([[S:[ ]*]][[NUM_TEAMS:[0-9]+]]X 128)
-/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:6 ConstWGSize:64  args: 6 teamsXthrds:([[S:[ ]*]][[NUM_TEAMS:[0-9]+]]X 64)
-/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:6 ConstWGSize:256  args: 6 teamsXthrds:([[S:[ ]*]][[NUM_TEAMS:[0-9]+]]X  32)
+/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:6 ConstWGSize:128  args: 6 teamsXthrds:({{ *}}1X 128)
+/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:6 ConstWGSize:64  args: 6 teamsXthrds:({{ *}}1X  64)
+/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:6 ConstWGSize:256  args: 6 teamsXthrds:({{ *}}1X 256)
 
diff --git a/test/smoke-dev/no-loop-split-3/no_loop_split_3.c b/test/smoke-dev/no-loop-split-3/no_loop_split_3.c
index 7df3433e4b..494e9d68ef 100644
--- a/test/smoke-dev/no-loop-split-3/no_loop_split_3.c
+++ b/test/smoke-dev/no-loop-split-3/no_loop_split_3.c
@@ -97,13 +97,15 @@ int main()
   return rc;
 }
 
-/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:6 ConstWGSize:1024  args: 6 teamsXthrds:([[S:[ ]*]][[NUM_TEAMS:[0-9]+]]X  64)
-/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:6 ConstWGSize:1024  args: 6 teamsXthrds:([[S:[ ]*]][[NUM_TEAMS:[0-9]+]]X 100)
-/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:6 ConstWGSize:1024  args: 6 teamsXthrds:([[S:[ ]*]][[NUM_TEAMS:[0-9]+]]X  32)
-/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:6 ConstWGSize:1024  args: 6 teamsXthrds:([[S:[ ]*]][[NUM_TEAMS:[0-9]+]]X  32)
-/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:6 ConstWGSize:1024  args: 6 teamsXthrds:([[S:[ ]*]][[NUM_TEAMS:[0-9]+]]X  32)
+// Every loop here is small. Its kernel gets the full block size (ConstWGSize,
+// or the thread_limit clause) and as many teams as its trip count needs.
+/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:6 ConstWGSize:1024  args: 6 teamsXthrds:({{ *}}1X  64)
+/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:6 ConstWGSize:1024  args: 6 teamsXthrds:({{ *}}1X 100)
+/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:6 ConstWGSize:1024  args: 6 teamsXthrds:({{ *}}1X1024)
+/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:6 ConstWGSize:1024  args: 6 teamsXthrds:({{ *}}1X1024)
+/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:6 ConstWGSize:1024  args: 6 teamsXthrds:({{ *}}1X1024)
 /// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:3 ConstWGSize:{{288|320}}  args: 8 teamsXthrds:([[S:[ ]*]][[NUM_TEAMS:[0-9]+]]X 256)
-/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:6 ConstWGSize:1024  args: 6 teamsXthrds:([[S:[ ]*]][[NUM_TEAMS:[0-9]+]]X 128)
-/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:6 ConstWGSize:1024  args: 6 teamsXthrds:([[S:[ ]*]][[NUM_TEAMS:[0-9]+]]X 64)
-/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:6 ConstWGSize:1024  args: 6 teamsXthrds:([[S:[ ]*]][[NUM_TEAMS:[0-9]+]]X  32)
+/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:6 ConstWGSize:1024  args: 6 teamsXthrds:({{ *}}1X 128)
+/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:6 ConstWGSize:1024  args: 6 teamsXthrds:({{ *}}1X  64)
+/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:6 ConstWGSize:1024  args: 6 teamsXthrds:({{ *}}1X1024)
 
diff --git a/test/smoke-dev/occupancy-based-opt-generic-spmd-use-small-blocksize/Makefile b/test/smoke-dev/occupancy-based-opt-generic-spmd-use-small-blocksize/Makefile
deleted file mode 100644
index abfdd20c7b..0000000000
--- a/test/smoke-dev/occupancy-based-opt-generic-spmd-use-small-blocksize/Makefile
+++ /dev/null
@@ -1,18 +0,0 @@
-include ../../Makefile.defs
-
-TESTNAME     = occupancy-based-opt-generic-spmd-use-small-blocksize
-TESTSRC_MAIN = occupancy-based-opt-generic-spmd-use-small-blocksize.c
-TESTSRC_AUX  =
-TESTSRC_ALL  = $(TESTSRC_MAIN) $(TESTSRC_AUX)
-RUNENV      += LIBOMPTARGET_KERNEL_TRACE=1
-RUNENV 	    += OMPX_GENERIC_SPMD_OCCUPANCY_BASED_OPT=1
-
-RUNCMD      = ./$(TESTNAME) 2>&1 | $(FILECHECK) $(TESTSRC_MAIN)
-
-CLANG        ?= clang
-OMP_BIN      = $(AOMP)/bin/$(CLANG)
-CC           = $(OMP_BIN) $(VERBOSE)
-#-ccc-print-phases
-#"-\#\#\#"
-
-include ../Makefile.rules
diff --git a/test/smoke-dev/occupancy-based-opt-generic-spmd-use-small-blocksize/occupancy-based-opt-generic-spmd-use-small-blocksize.c b/test/smoke-dev/occupancy-based-opt-generic-spmd-use-small-blocksize/occupancy-based-opt-generic-spmd-use-small-blocksize.c
deleted file mode 100644
index c10b93f92a..0000000000
--- a/test/smoke-dev/occupancy-based-opt-generic-spmd-use-small-blocksize/occupancy-based-opt-generic-spmd-use-small-blocksize.c
+++ /dev/null
@@ -1,36 +0,0 @@
-// Copyright © Advanced Micro Devices, Inc., or its affiliates.
-//
-// SPDX-License-Identifier:  MIT
-
-#define N 4194304
-#define NUM_THREADS 256
-#define NUM_DISTR (N / NUM_THREADS)
-
-#include <stdio.h>
-#include <stdlib.h>
-
-int main() {
-  int *a = (int*)malloc(sizeof(int)*N);
-  int i, j;
-
-#pragma omp target teams distribute map(tofrom:a[0:N])
-  for (i = 0; i < NUM_DISTR; i++)
-#pragma omp parallel for
-    for (j = 0; j < NUM_THREADS; j++)
-      a[i * NUM_THREADS + j] = (i * NUM_THREADS + j);
-
-  for (i = 0; i < N; i++)
-    if (a[i] != i) {
-      printf("wrong value: a[%d]=%d\n", i, a[i]);
-      free(a);
-      return 1;
-    }
-  printf("Success\n");
-  free(a);
-  return 0;
-}
-
-/// CHECK: SGN:3
-/// CHECK: teamsXthrds:([[S:[ ]*]][[NUM_TEAMS:[0-9]+]]X  64)
-/// CHECK: Achieved Occupancy: 100%
-
diff --git a/test/smoke-dev/occupancy-based-opt-generic-spmd/Makefile b/test/smoke-dev/occupancy-based-opt-generic-spmd/Makefile
deleted file mode 100644
index 08e02b189d..0000000000
--- a/test/smoke-dev/occupancy-based-opt-generic-spmd/Makefile
+++ /dev/null
@@ -1,18 +0,0 @@
-include ../../Makefile.defs
-
-TESTNAME     = occupancy_based_opt_generic_spmd
-TESTSRC_MAIN = occupancy_based_opt_generic_spmd.c
-TESTSRC_AUX  =
-TESTSRC_ALL  = $(TESTSRC_MAIN) $(TESTSRC_AUX)
-RUNENV      += LIBOMPTARGET_KERNEL_TRACE=1
-RUNENV 	    += OMPX_GENERIC_SPMD_OCCUPANCY_BASED_OPT=1
-
-RUNCMD      = ./$(TESTNAME) 2>&1 | $(FILECHECK) $(TESTSRC_MAIN)
-
-CLANG        ?= clang
-OMP_BIN      = $(AOMP)/bin/$(CLANG)
-CC           = $(OMP_BIN) $(VERBOSE)
-#-ccc-print-phases
-#"-\#\#\#"
-
-include ../Makefile.rules
diff --git a/test/smoke-dev/occupancy-based-opt-generic-spmd/occupancy_based_opt_generic_spmd.c b/test/smoke-dev/occupancy-based-opt-generic-spmd/occupancy_based_opt_generic_spmd.c
deleted file mode 100644
index 3a1a2bb135..0000000000
--- a/test/smoke-dev/occupancy-based-opt-generic-spmd/occupancy_based_opt_generic_spmd.c
+++ /dev/null
@@ -1,30 +0,0 @@
-// Copyright © Advanced Micro Devices, Inc., or its affiliates.
-//
-// SPDX-License-Identifier:  MIT
-
-#define N 1048576
-#define NUM_THREADS 256
-#define NUM_DISTR (N / NUM_THREADS)
-
-#include <stdio.h>
-
-int main() {
-  int a[N];
-  int i, j;
-
-#pragma omp target teams distribute
-  for (i = 0; i < NUM_DISTR; i++)
-#pragma omp parallel for
-    for (j = 0; j < NUM_THREADS; j++)
-      a[i * NUM_THREADS + j] = (i * NUM_THREADS + j);
-
-  for (i = 0; i < N; i++)
-    if (a[i] != i) {
-      printf("wrong value: a[%d]=%d\n", i, a[i]);
-      return 1;
-    }
-  printf("Success\n");
-  return 0;
-}
-
-/// CHECK: Achieved Occupancy: 100%
diff --git a/test/smoke-dev/occupancy-based-opt-xteam-reduction/Makefile b/test/smoke-dev/occupancy-based-opt-xteam-reduction/Makefile
deleted file mode 100644
index 6be7601355..0000000000
--- a/test/smoke-dev/occupancy-based-opt-xteam-reduction/Makefile
+++ /dev/null
@@ -1,18 +0,0 @@
-include ../../Makefile.defs
-
-TESTNAME     = occupancy_based_opt_xteam_reduction
-TESTSRC_MAIN = occupancy_based_opt_xteam_reduction.c
-TESTSRC_AUX  =
-TESTSRC_ALL  = $(TESTSRC_MAIN) $(TESTSRC_AUX)
-RUNENV      += LIBOMPTARGET_KERNEL_TRACE=1 LIBOMPTARGET_AMDGPU_XTEAM_BLOCKSIZE=512
-RUNENV      += OMPX_XTEAMREDUCTION_OCCUPANCY_BASED_OPT=1
-
-RUNCMD      = ./$(TESTNAME) 2>&1 | $(FILECHECK) $(TESTSRC_MAIN)
-
-CLANG        ?= clang
-OMP_BIN      = $(AOMP)/bin/$(CLANG)
-CC           = $(OMP_BIN) $(VERBOSE)
-#-ccc-print-phases
-#"-\#\#\#"
-
-include ../Makefile.rules
diff --git a/test/smoke-dev/occupancy-based-opt-xteam-reduction/occupancy_based_opt_xteam_reduction.c b/test/smoke-dev/occupancy-based-opt-xteam-reduction/occupancy_based_opt_xteam_reduction.c
deleted file mode 100644
index 6e548056a4..0000000000
--- a/test/smoke-dev/occupancy-based-opt-xteam-reduction/occupancy_based_opt_xteam_reduction.c
+++ /dev/null
@@ -1,51 +0,0 @@
-// Copyright © Advanced Micro Devices, Inc., or its affiliates.
-//
-// SPDX-License-Identifier:  MIT
-
-#include <stdio.h>
-#include <omp.h>
-
-int main()
-{
-  int N = 1000000;
-
-  double a[N];
-
-  for (int i=0; i<N; i++)
-    a[i]=i;
-
-  double sum1, sum2, sum3, sum4;
-  sum1 = sum2 = sum3 = sum4 = 0;
-  
-#pragma omp target teams distribute parallel for map(tofrom:sum1) reduction(+:sum1)
-  for (int j = 0; j< N; j=j+1)
-    sum1 += a[j];
-
-#pragma omp target teams distribute parallel for map(tofrom:sum2) reduction(+:sum2)
-  for (int j = 0; j< N; j=j+2)
-    sum2 += a[j];
-
-#pragma omp target teams distribute parallel for map(tofrom:sum3,sum4) reduction(+:sum3,sum4)
-  for (int j = 0; j< N; j=j+1) {
-    sum3 += a[j];
-    sum4 += a[j];
-  }
-  
-  printf("%f %f %f %f\n", sum1, sum2, sum3, sum4);
-  
-  int rc =
-    (sum1 != 499999500000) ||
-    (sum2 != 249999500000) ||
-    (sum3 != 499999500000) ||
-    (sum4 != 499999500000);
-
-  if (!rc)
-    printf("Success\n");
-
-  return rc;
-}
-
-/// CHECK: Achieved Occupancy: 100%
-/// CHECK: Achieved Occupancy: 100%
-/// CHECK: Achieved Occupancy: 100%
-
diff --git a/test/smoke-dev/xteam-red-cmdline-default-option/Makefile b/test/smoke-dev/xteam-red-cmdline-default-option/Makefile
index b502b35072..d883adafef 100644
--- a/test/smoke-dev/xteam-red-cmdline-default-option/Makefile
+++ b/test/smoke-dev/xteam-red-cmdline-default-option/Makefile
@@ -4,7 +4,7 @@ TESTNAME     = xteam_red_cmdline_option1
 TESTSRC_MAIN = xteam_red_cmdline_option1.c
 TESTSRC_AUX  =
 TESTSRC_ALL  = $(TESTSRC_MAIN) $(TESTSRC_AUX)
-RUNENV      += LIBOMPTARGET_KERNEL_TRACE=1 LIBOMPTARGET_AMDGPU_XTEAM_BLOCKSIZE=512
+RUNENV      += LIBOMPTARGET_KERNEL_TRACE=1
 
 RUNCMD      = ./$(TESTNAME) 2>&1 | $(FILECHECK) $(TESTSRC_MAIN)
 
diff --git a/test/smoke-dev/xteam-red-cmdline-option1/Makefile b/test/smoke-dev/xteam-red-cmdline-option1/Makefile
index 3ec91db265..8498a7231b 100644
--- a/test/smoke-dev/xteam-red-cmdline-option1/Makefile
+++ b/test/smoke-dev/xteam-red-cmdline-option1/Makefile
@@ -4,7 +4,7 @@ TESTNAME     = xteam_red_cmdline_option1
 TESTSRC_MAIN = xteam_red_cmdline_option1.c
 TESTSRC_AUX  =
 TESTSRC_ALL  = $(TESTSRC_MAIN) $(TESTSRC_AUX)
-RUNENV      += LIBOMPTARGET_KERNEL_TRACE=1 LIBOMPTARGET_AMDGPU_XTEAM_BLOCKSIZE=512
+RUNENV      += LIBOMPTARGET_KERNEL_TRACE=1
 
 RUNCMD      = ./$(TESTNAME) 2>&1 | $(FILECHECK) $(TESTSRC_MAIN)
 
diff --git a/test/smoke-dev/xteam-red-split-1/Makefile b/test/smoke-dev/xteam-red-split-1/Makefile
index d268953eb7..a84cd95275 100644
--- a/test/smoke-dev/xteam-red-split-1/Makefile
+++ b/test/smoke-dev/xteam-red-split-1/Makefile
@@ -4,7 +4,7 @@ TESTNAME     = xteam_red_split_1
 TESTSRC_MAIN = xteam_red_split_1.c
 TESTSRC_AUX  =
 TESTSRC_ALL  = $(TESTSRC_MAIN) $(TESTSRC_AUX)
-RUNENV      += LIBOMPTARGET_KERNEL_TRACE=1 LIBOMPTARGET_AMDGPU_XTEAM_BLOCKSIZE=512
+RUNENV      += LIBOMPTARGET_KERNEL_TRACE=1
 
 RUNCMD      = ./$(TESTNAME) 2>&1 | $(FILECHECK) $(TESTSRC_MAIN)
 
diff --git a/test/smoke-dev/xteam-red-split-1/xteam_red_split_1.c b/test/smoke-dev/xteam-red-split-1/xteam_red_split_1.c
index 06943f1893..51dc5a148f 100644
--- a/test/smoke-dev/xteam-red-split-1/xteam_red_split_1.c
+++ b/test/smoke-dev/xteam-red-split-1/xteam_red_split_1.c
@@ -150,17 +150,20 @@ int main()
 // required: without it every team accumulates into the same shared variable,
 // which is a data race. The removed downstream implementation used to paper
 // over that by always reducing across teams.
-/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:512 args: 5 teamsXthrds:([[S:[ ]*]][[NUM_TEAMS:[0-9]+]]X 32)
-/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:64 args: 5 teamsXthrds:([[S:[ ]*]][[NUM_TEAMS:[0-9]+]]X 64)
-/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:512 args: 5 teamsXthrds:([[S:[ ]*]][[NUM_TEAMS:[0-9]+]]X 32)
-/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:128 args: 5 teamsXthrds:([[S:[ ]*]][[NUM_TEAMS:[0-9]+]]X 32)
-/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:512 args: 5 teamsXthrds:([[S:[ ]*]][[NUM_TEAMS:[0-9]+]]X 32)
-/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:128 args: 5 teamsXthrds:([[S:[ ]*]][[NUM_TEAMS:[0-9]+]]X 32)
-/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:512 args: 5 teamsXthrds:([[S:[ ]*]][[NUM_TEAMS:[0-9]+]]X 32)
-/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:128 args: 5 teamsXthrds:([[S:[ ]*]][[NUM_TEAMS:[0-9]+]]X 128)
-/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:512 args: 5 teamsXthrds:([[S:[ ]*]][[NUM_TEAMS:[0-9]+]]X 32)
-/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:256 args: 5 teamsXthrds:([[S:[ ]*]][[NUM_TEAMS:[0-9]+]]X 256)
-/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:512 args: 5 teamsXthrds:([[S:[ ]*]][[NUM_TEAMS:[0-9]+]]X 32)
-/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:128 args: 5 teamsXthrds:([[S:[ ]*]][[NUM_TEAMS:[0-9]+]]X 128)
-/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:512 args: 6 teamsXthrds:([[S:[ ]*]][[NUM_TEAMS:[0-9]+]]X 32)
-/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:64 args: 6 teamsXthrds:([[S:[ ]*]][[NUM_TEAMS:[0-9]+]]X 64)
+//
+// Every loop here is small. Its kernel gets the full block size (ConstWGSize,
+// or the thread_limit clause) and as many teams as its trip count needs.
+/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:512 args: 5 teamsXthrds:({{ *}}2X 512)
+/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:64 args: 5 teamsXthrds:({{ *}}16X  64)
+/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:512 args: 5 teamsXthrds:({{ *}}2X 512)
+/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:128 args: 5 teamsXthrds:({{ *}}8X 128)
+/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:512 args: 5 teamsXthrds:({{ *}}2X 512)
+/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:128 args: 5 teamsXthrds:({{ *}}8X 128)
+/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:512 args: 5 teamsXthrds:({{ *}}2X 512)
+/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:128 args: 5 teamsXthrds:({{ *}}8X 128)
+/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:512 args: 5 teamsXthrds:({{ *}}2X 512)
+/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:256 args: 5 teamsXthrds:({{ *}}4X 256)
+/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:512 args: 5 teamsXthrds:({{ *}}1X 512)
+/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:128 args: 5 teamsXthrds:({{ *}}4X 128)
+/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:512 args: 6 teamsXthrds:({{ *}}2X 512)
+/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:64 args: 6 teamsXthrds:({{ *}}16X  64)
diff --git a/test/smoke-dev/xteam-red-split-2/xteam_red_split_2.c b/test/smoke-dev/xteam-red-split-2/xteam_red_split_2.c
index a4665c53c4..fc4ffeaca4 100644
--- a/test/smoke-dev/xteam-red-split-2/xteam_red_split_2.c
+++ b/test/smoke-dev/xteam-red-split-2/xteam_red_split_2.c
@@ -156,17 +156,20 @@ int main()
 // '-fopenmp-target-xteam-reduction-blocksize=' overrides the thread_limit and
 // num_threads clauses on the constructs below, which is how the block size of
 // a cross-team reduction kernel has always been selected.
-/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:64 args: 5 teamsXthrds:([[S:[ ]*]][[NUM_TEAMS:[0-9]+]]X 32)
-/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:64 args: 5 teamsXthrds:([[S:[ ]*]][[NUM_TEAMS:[0-9]+]]X 64)
-/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:64 args: 5 teamsXthrds:([[S:[ ]*]][[NUM_TEAMS:[0-9]+]]X 32)
-/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:64 args: 5 teamsXthrds:([[S:[ ]*]][[NUM_TEAMS:[0-9]+]]X 32)
-/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:64 args: 5 teamsXthrds:([[S:[ ]*]][[NUM_TEAMS:[0-9]+]]X 32)
-/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:64 args: 5 teamsXthrds:([[S:[ ]*]][[NUM_TEAMS:[0-9]+]]X 32)
-/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:64 args: 5 teamsXthrds:([[S:[ ]*]][[NUM_TEAMS:[0-9]+]]X 32)
-/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:64 args: 5 teamsXthrds:([[S:[ ]*]][[NUM_TEAMS:[0-9]+]]X 64)
-/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:64 args: 5 teamsXthrds:([[S:[ ]*]][[NUM_TEAMS:[0-9]+]]X 32)
-/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:64 args: 5 teamsXthrds:([[S:[ ]*]][[NUM_TEAMS:[0-9]+]]X 64)
-/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:64 args: 5 teamsXthrds:([[S:[ ]*]][[NUM_TEAMS:[0-9]+]]X 32)
-/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:64 args: 5 teamsXthrds:([[S:[ ]*]][[NUM_TEAMS:[0-9]+]]X 64)
-/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:64 args: 6 teamsXthrds:([[S:[ ]*]][[NUM_TEAMS:[0-9]+]]X 32)
-/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:64 args: 6 teamsXthrds:([[S:[ ]*]][[NUM_TEAMS:[0-9]+]]X 64)
+//
+// Every loop here is small. Its kernel still gets the full block size and as
+// many teams as its trip count needs.
+/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:64 args: 5 teamsXthrds:({{ *}}16X  64)
+/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:64 args: 5 teamsXthrds:({{ *}}16X  64)
+/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:64 args: 5 teamsXthrds:({{ *}}16X  64)
+/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:64 args: 5 teamsXthrds:({{ *}}16X  64)
+/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:64 args: 5 teamsXthrds:({{ *}}16X  64)
+/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:64 args: 5 teamsXthrds:({{ *}}16X  64)
+/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:64 args: 5 teamsXthrds:({{ *}}16X  64)
+/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:64 args: 5 teamsXthrds:({{ *}}16X  64)
+/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:64 args: 5 teamsXthrds:({{ *}}16X  64)
+/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:64 args: 5 teamsXthrds:({{ *}}16X  64)
+/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:64 args: 5 teamsXthrds:({{ *}}8X  64)
+/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:64 args: 5 teamsXthrds:({{ *}}8X  64)
+/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:64 args: 6 teamsXthrds:({{ *}}16X  64)
+/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:64 args: 6 teamsXthrds:({{ *}}16X  64)
diff --git a/test/smoke-dev/xteam-red-split-default-option-1/Makefile b/test/smoke-dev/xteam-red-split-default-option-1/Makefile
index e4b3aecf6e..eb2aabb7c5 100644
--- a/test/smoke-dev/xteam-red-split-default-option-1/Makefile
+++ b/test/smoke-dev/xteam-red-split-default-option-1/Makefile
@@ -4,7 +4,7 @@ TESTNAME     = xteam_red_split_1
 TESTSRC_MAIN = xteam_red_split_1.c
 TESTSRC_AUX  =
 TESTSRC_ALL  = $(TESTSRC_MAIN) $(TESTSRC_AUX)
-RUNENV      += LIBOMPTARGET_KERNEL_TRACE=1 LIBOMPTARGET_AMDGPU_XTEAM_BLOCKSIZE=512
+RUNENV      += LIBOMPTARGET_KERNEL_TRACE=1
 
 RUNCMD      = ./$(TESTNAME) 2>&1 | $(FILECHECK) $(TESTSRC_MAIN)
 
diff --git a/test/smoke-dev/xteam-red-split-default-option-1/xteam_red_split_1.c b/test/smoke-dev/xteam-red-split-default-option-1/xteam_red_split_1.c
index 06943f1893..51dc5a148f 100644
--- a/test/smoke-dev/xteam-red-split-default-option-1/xteam_red_split_1.c
+++ b/test/smoke-dev/xteam-red-split-default-option-1/xteam_red_split_1.c
@@ -150,17 +150,20 @@ int main()
 // required: without it every team accumulates into the same shared variable,
 // which is a data race. The removed downstream implementation used to paper
 // over that by always reducing across teams.
-/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:512 args: 5 teamsXthrds:([[S:[ ]*]][[NUM_TEAMS:[0-9]+]]X 32)
-/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:64 args: 5 teamsXthrds:([[S:[ ]*]][[NUM_TEAMS:[0-9]+]]X 64)
-/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:512 args: 5 teamsXthrds:([[S:[ ]*]][[NUM_TEAMS:[0-9]+]]X 32)
-/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:128 args: 5 teamsXthrds:([[S:[ ]*]][[NUM_TEAMS:[0-9]+]]X 32)
-/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:512 args: 5 teamsXthrds:([[S:[ ]*]][[NUM_TEAMS:[0-9]+]]X 32)
-/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:128 args: 5 teamsXthrds:([[S:[ ]*]][[NUM_TEAMS:[0-9]+]]X 32)
-/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:512 args: 5 teamsXthrds:([[S:[ ]*]][[NUM_TEAMS:[0-9]+]]X 32)
-/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:128 args: 5 teamsXthrds:([[S:[ ]*]][[NUM_TEAMS:[0-9]+]]X 128)
-/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:512 args: 5 teamsXthrds:([[S:[ ]*]][[NUM_TEAMS:[0-9]+]]X 32)
-/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:256 args: 5 teamsXthrds:([[S:[ ]*]][[NUM_TEAMS:[0-9]+]]X 256)
-/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:512 args: 5 teamsXthrds:([[S:[ ]*]][[NUM_TEAMS:[0-9]+]]X 32)
-/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:128 args: 5 teamsXthrds:([[S:[ ]*]][[NUM_TEAMS:[0-9]+]]X 128)
-/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:512 args: 6 teamsXthrds:([[S:[ ]*]][[NUM_TEAMS:[0-9]+]]X 32)
-/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:64 args: 6 teamsXthrds:([[S:[ ]*]][[NUM_TEAMS:[0-9]+]]X 64)
+//
+// Every loop here is small. Its kernel gets the full block size (ConstWGSize,
+// or the thread_limit clause) and as many teams as its trip count needs.
+/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:512 args: 5 teamsXthrds:({{ *}}2X 512)
+/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:64 args: 5 teamsXthrds:({{ *}}16X  64)
+/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:512 args: 5 teamsXthrds:({{ *}}2X 512)
+/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:128 args: 5 teamsXthrds:({{ *}}8X 128)
+/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:512 args: 5 teamsXthrds:({{ *}}2X 512)
+/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:128 args: 5 teamsXthrds:({{ *}}8X 128)
+/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:512 args: 5 teamsXthrds:({{ *}}2X 512)
+/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:128 args: 5 teamsXthrds:({{ *}}8X 128)
+/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:512 args: 5 teamsXthrds:({{ *}}2X 512)
+/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:256 args: 5 teamsXthrds:({{ *}}4X 256)
+/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:512 args: 5 teamsXthrds:({{ *}}1X 512)
+/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:128 args: 5 teamsXthrds:({{ *}}4X 128)
+/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:512 args: 6 teamsXthrds:({{ *}}2X 512)
+/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:64 args: 6 teamsXthrds:({{ *}}16X  64)
diff --git a/test/smoke-dev/xteam-red-split-default-option-2/xteam_red_split_2.c b/test/smoke-dev/xteam-red-split-default-option-2/xteam_red_split_2.c
index 58ed88be65..d1147ff501 100644
--- a/test/smoke-dev/xteam-red-split-default-option-2/xteam_red_split_2.c
+++ b/test/smoke-dev/xteam-red-split-default-option-2/xteam_red_split_2.c
@@ -151,17 +151,20 @@ int main()
 // '-fopenmp-target-xteam-reduction-blocksize=' overrides the thread_limit and
 // num_threads clauses on the constructs below, which is how the block size of
 // a cross-team reduction kernel has always been selected.
-/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:64 args: 5 teamsXthrds:([[S:[ ]*]][[NUM_TEAMS:[0-9]+]]X 32)
-/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:64 args: 5 teamsXthrds:([[S:[ ]*]][[NUM_TEAMS:[0-9]+]]X 64)
-/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:64 args: 5 teamsXthrds:([[S:[ ]*]][[NUM_TEAMS:[0-9]+]]X 32)
-/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:64 args: 5 teamsXthrds:([[S:[ ]*]][[NUM_TEAMS:[0-9]+]]X 32)
-/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:64 args: 5 teamsXthrds:([[S:[ ]*]][[NUM_TEAMS:[0-9]+]]X 32)
-/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:64 args: 5 teamsXthrds:([[S:[ ]*]][[NUM_TEAMS:[0-9]+]]X 32)
-/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:64 args: 5 teamsXthrds:([[S:[ ]*]][[NUM_TEAMS:[0-9]+]]X 32)
-/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:64 args: 5 teamsXthrds:([[S:[ ]*]][[NUM_TEAMS:[0-9]+]]X 64)
-/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:64 args: 5 teamsXthrds:([[S:[ ]*]][[NUM_TEAMS:[0-9]+]]X 32)
-/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:64 args: 5 teamsXthrds:([[S:[ ]*]][[NUM_TEAMS:[0-9]+]]X 64)
-/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:64 args: 5 teamsXthrds:([[S:[ ]*]][[NUM_TEAMS:[0-9]+]]X 32)
-/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:64 args: 5 teamsXthrds:([[S:[ ]*]][[NUM_TEAMS:[0-9]+]]X 64)
-/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:64 args: 6 teamsXthrds:([[S:[ ]*]][[NUM_TEAMS:[0-9]+]]X 32)
-/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:64 args: 6 teamsXthrds:([[S:[ ]*]][[NUM_TEAMS:[0-9]+]]X 64)
+//
+// Every loop here is small. Its kernel still gets the full block size and as
+// many teams as its trip count needs.
+/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:64 args: 5 teamsXthrds:({{ *}}16X  64)
+/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:64 args: 5 teamsXthrds:({{ *}}16X  64)
+/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:64 args: 5 teamsXthrds:({{ *}}16X  64)
+/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:64 args: 5 teamsXthrds:({{ *}}16X  64)
+/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:64 args: 5 teamsXthrds:({{ *}}16X  64)
+/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:64 args: 5 teamsXthrds:({{ *}}16X  64)
+/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:64 args: 5 teamsXthrds:({{ *}}16X  64)
+/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:64 args: 5 teamsXthrds:({{ *}}16X  64)
+/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:64 args: 5 teamsXthrds:({{ *}}16X  64)
+/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:64 args: 5 teamsXthrds:({{ *}}16X  64)
+/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:64 args: 5 teamsXthrds:({{ *}}8X  64)
+/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:64 args: 5 teamsXthrds:({{ *}}8X  64)
+/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:64 args: 6 teamsXthrds:({{ *}}16X  64)
+/// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 ConstWGSize:64 args: 6 teamsXthrds:({{ *}}16X  64)
diff --git a/test/smoke-dev/xteam-red-teams-default-option/Makefile b/test/smoke-dev/xteam-red-teams-default-option/Makefile
index 8f8c18eda7..8653b9f8ff 100644
--- a/test/smoke-dev/xteam-red-teams-default-option/Makefile
+++ b/test/smoke-dev/xteam-red-teams-default-option/Makefile
@@ -4,7 +4,7 @@ TESTNAME     = xteam_red_teams
 TESTSRC_MAIN = xteam_red_teams.c
 TESTSRC_AUX  =
 TESTSRC_ALL  = $(TESTSRC_MAIN) $(TESTSRC_AUX)
-RUNENV      += LIBOMPTARGET_KERNEL_TRACE=1 LIBOMPTARGET_AMDGPU_XTEAM_BLOCKSIZE=512
+RUNENV      += LIBOMPTARGET_KERNEL_TRACE=1
 
 RUNCMD      = ./$(TESTNAME) 2>&1 | $(FILECHECK) $(TESTSRC_MAIN)
 
diff --git a/test/smoke-dev/xteam-red-teams/Makefile b/test/smoke-dev/xteam-red-teams/Makefile
index aedd0cf3d5..03f3eb79c7 100644
--- a/test/smoke-dev/xteam-red-teams/Makefile
+++ b/test/smoke-dev/xteam-red-teams/Makefile
@@ -4,7 +4,7 @@ TESTNAME     = xteam_red_teams
 TESTSRC_MAIN = xteam_red_teams.c
 TESTSRC_AUX  =
 TESTSRC_ALL  = $(TESTSRC_MAIN) $(TESTSRC_AUX)
-RUNENV      += LIBOMPTARGET_KERNEL_TRACE=1 LIBOMPTARGET_AMDGPU_XTEAM_BLOCKSIZE=512
+RUNENV      += LIBOMPTARGET_KERNEL_TRACE=1
 
 RUNCMD      = ./$(TESTNAME) 2>&1 | $(FILECHECK) $(TESTSRC_MAIN)
 
diff --git a/test/smoke-limbo/occupancy-based-opt-big-jump-loop/Makefile b/test/smoke-limbo/occupancy-based-opt-big-jump-loop/Makefile
deleted file mode 100644
index 3281238c21..0000000000
--- a/test/smoke-limbo/occupancy-based-opt-big-jump-loop/Makefile
+++ /dev/null
@@ -1,18 +0,0 @@
-include ../../Makefile.defs
-
-TESTNAME     = occupancy_based_opt_big_jump_loop
-TESTSRC_MAIN = occupancy_based_opt_big_jump_loop.c
-TESTSRC_AUX  =
-TESTSRC_ALL  = $(TESTSRC_MAIN) $(TESTSRC_AUX)
-RUNENV      += LIBOMPTARGET_KERNEL_TRACE=1
-RUNENV 		+= OMPX_BIGJUMPLOOP_OCCUPANCY_BASED_OPT=1
-
-RUNCMD      = ./$(TESTNAME) 2>&1 | $(FILECHECK) $(TESTSRC_MAIN)
-
-CLANG        ?= clang
-OMP_BIN      = $(AOMP)/bin/$(CLANG)
-CC           = $(OMP_BIN) $(VERBOSE)
-#-ccc-print-phases
-#"-\#\#\#"
-
-include ../Makefile.rules
diff --git a/test/smoke-limbo/occupancy-based-opt-big-jump-loop/occupancy_based_opt_big_jump_loop.c b/test/smoke-limbo/occupancy-based-opt-big-jump-loop/occupancy_based_opt_big_jump_loop.c
deleted file mode 100755
index 26e37d921a..0000000000
--- a/test/smoke-limbo/occupancy-based-opt-big-jump-loop/occupancy_based_opt_big_jump_loop.c
+++ /dev/null
@@ -1,41 +0,0 @@
-// Copyright © Advanced Micro Devices, Inc., or its affiliates.
-//
-// SPDX-License-Identifier:  MIT
-
-#include <omp.h>
-#include <stdio.h>
-
-int main() {
-  int N = 1000000;
-
-  int a[N];
-  int b[N];
-
-  int i;
-
-  for (i = 0; i < N; i++)
-    a[i] = 0;
-
-  for (i = 0; i < N; i++)
-    b[i] = i;
-
-#pragma omp target teams distribute parallel for
-  {
-    for (int j = 0; j < N; j++)
-      a[j] = b[j];
-  }
-
-  int rc = 0;
-  for (i = 0; i < N; i++)
-    if (a[i] != b[i]) {
-      rc++;
-      printf("Wrong varlue: a[%d]=%d\n", i, a[i]);
-    }
-
-  if (!rc)
-    printf("Success\n");
-
-  return rc;
-}
-
-/// CHECK: Achieved Occupancy: 100%
diff --git a/test/smoke-limbo/occupancy-based-opt/Makefile b/test/smoke-limbo/occupancy-based-opt/Makefile
deleted file mode 100644
index aa50761a73..0000000000
--- a/test/smoke-limbo/occupancy-based-opt/Makefile
+++ /dev/null
@@ -1,19 +0,0 @@
-include ../../Makefile.defs
-
-TESTNAME     = occupancy_based_opt
-TESTSRC_MAIN = occupancy_based_opt.c
-TESTSRC_AUX  =
-TESTSRC_ALL  = $(TESTSRC_MAIN) $(TESTSRC_AUX)
-RUNENV      += LIBOMPTARGET_KERNEL_TRACE=1
-RUNENV 		+= OMPX_SPMD_OCCUPANCY_BASED_OPT=1
-CFLAGS      += -fno-openmp-target-big-jump-loop
-
-RUNCMD      = ./$(TESTNAME) 2>&1 | $(FILECHECK) $(TESTSRC_MAIN)
-
-CLANG        ?= clang
-OMP_BIN      = $(AOMP)/bin/$(CLANG)
-CC           = $(OMP_BIN) $(VERBOSE)
-#-ccc-print-phases
-#"-\#\#\#"
-
-include ../Makefile.rules
diff --git a/test/smoke-limbo/occupancy-based-opt/occupancy_based_opt.c b/test/smoke-limbo/occupancy-based-opt/occupancy_based_opt.c
deleted file mode 100644
index 38d9741580..0000000000
--- a/test/smoke-limbo/occupancy-based-opt/occupancy_based_opt.c
+++ /dev/null
@@ -1,42 +0,0 @@
-// Copyright © Advanced Micro Devices, Inc., or its affiliates.
-//
-// SPDX-License-Identifier:  MIT
-
-#include <stdio.h>
-#include <omp.h>
-
-int main()
-{
-  int N = 1000000;
-
-  int a[N];
-  int b[N];
-
-  int i;
-
-  for (i=0; i<N; i++)
-    a[i]=0;
-
-  for (i=0; i<N; i++)
-    b[i]=i;
-
-#pragma omp target teams distribute parallel for 
-  {
-    for (int j = 0; j< N; j++)
-      a[j]=b[j];
-  }
-
-  int rc = 0;
-  for (i=0; i<N; i++)
-    if (a[i] != b[i] ) {
-      rc++;
-      printf ("Wrong varlue: a[%d]=%d\n", i, a[i]);
-    }
-
-  if (!rc)
-    printf("Success\n");
-
-  return rc;
-}
-
-/// CHECK: Achieved Occupancy: 100%
diff --git a/test/smoke/stacked-block-size-config/Makefile b/test/smoke/stacked-block-size-config/Makefile
index 37fad41a27..da36275b9c 100644
--- a/test/smoke/stacked-block-size-config/Makefile
+++ b/test/smoke/stacked-block-size-config/Makefile
@@ -4,11 +4,12 @@ TESTNAME     = stacked_block_size_config
 TESTSRC_MAIN = stacked_block_size_config.c
 TESTSRC_AUX  =
 TESTSRC_ALL  = $(TESTSRC_MAIN) $(TESTSRC_AUX)
-RUNENV      += LIBOMPTARGET_MIN_THREADS_FOR_LOW_TRIP_COUNT=16 LIBOMPTARGET_AMDGPU_LOW_TRIPCOUNT=15001 LIBOMPTARGET_KERNEL_TRACE=1
+RUNENV      += LIBOMPTARGET_MIN_THREADS_FOR_LOW_TRIP_COUNT=16 LIBOMPTARGET_KERNEL_TRACE=1
 
 RUNCMD       = ./$(TESTNAME) 2>&1 | $(FILECHECK) $(TESTSRC_MAIN)
 
-CFLAGS       += -O3 -fopenmp-gpu-threads-per-team=128
+# The low trip count thread reduction applies to plain SPMD kernels.
+CFLAGS       += -O3 -fopenmp-gpu-threads-per-team=128 -fno-openmp-target-big-jump-loop
 CLANG        = clang
 OMP_BIN      = $(AOMP)/bin/$(CLANG)
 CC           = $(OMP_BIN) $(VERBOSE)
diff --git a/test/smoke/stacked-block-size-config/stacked_block_size_config.c b/test/smoke/stacked-block-size-config/stacked_block_size_config.c
index 52f4f6e27a..de14e8529a 100644
--- a/test/smoke/stacked-block-size-config/stacked_block_size_config.c
+++ b/test/smoke/stacked-block-size-config/stacked_block_size_config.c
@@ -1,7 +1,7 @@
 #include <stdio.h>
 #include <stdlib.h>
 
-#define SMALL 15000
+#define SMALL 1000
 #define LARGE 1000000
 
 int main() {
@@ -32,8 +32,10 @@ int main() {
   free(x);
   free(y);
 
-  /// CHECK: DEVID:{{.*}}SGN:7 ConstWGSize:128  args: 2 teamsXthrds:({{.*}}X  16) {{.*}}tripcount:15000 rpc:0
-  /// CHECK: DEVID:{{.*}}SGN:7 ConstWGSize:128  args: 2 teamsXthrds:({{.*}}X 128) {{.*}}tripcount:1000000 rpc:0
+  /// The block size from the command line is not a user override: the low
+  /// trip count rule still lowers the threads to the minimum.
+  /// CHECK: DEVID:{{.*}}SGN:2 ConstWGSize:128  args: 2 teamsXthrds:({{ *}}63X  16) {{.*}}tripcount:1000 rpc:0
+  /// CHECK: DEVID:{{.*}}SGN:2 ConstWGSize:128  args: 2 teamsXthrds:({{.*}}X 128) {{.*}}tripcount:1000000 rpc:0
 
   /// CHECK: x[3] = 4
   /// CHECK: y[500000] = 8