diff --git a/libgomp/config/gcn/icv-device.c b/libgomp/config/gcn/icv-device.c
index fcfa0f3..26b2e7a 100644
--- a/libgomp/config/gcn/icv-device.c
+++ b/libgomp/config/gcn/icv-device.c
@@ -28,6 +28,16 @@
 
 #include "libgomp.h"
 
+/* This is set to the ICV values of current GPU during device initialization,
+   when the offload image containing this libgomp portion is loaded.  */
+volatile int GOMP_DEVICE_NUM_VAR;
+volatile int GOMP_RUN_SCHED_VAR;
+volatile int GOMP_RUN_SCHED_CHUNK_SIZE;
+volatile int GOMP_DEFAULT_DEVICE_VAR;
+volatile int GOMP_MAX_ACTIVE_LEVELS_VAR;
+volatile omp_proc_bind_t GOMP_BIND_VAR;
+volatile int GOMP_NTEAMS_VAR;
+
 void
 omp_set_default_device (int device_num __attribute__((unused)))
 {
@@ -36,7 +46,7 @@ omp_set_default_device (int device_num __attribute__((unused)))
 int
 omp_get_default_device (void)
 {
-  return 0;
+  return GOMP_DEFAULT_DEVICE_VAR;
 }
 
 int
@@ -58,19 +68,22 @@ omp_is_initial_device (void)
   return 0;
 }
 
-/* This is set to the device number of current GPU during device initialization,
-   when the offload image containing this libgomp portion is loaded.  */
-static volatile int GOMP_DEVICE_NUM_VAR;
-
 int
 omp_get_device_num (void)
 {
   return GOMP_DEVICE_NUM_VAR;
 }
 
+int
+omp_get_max_teams (void)
+{
+  return GOMP_NTEAMS_VAR;
+}
+
 ialias (omp_set_default_device)
 ialias (omp_get_default_device)
 ialias (omp_get_initial_device)
 ialias (omp_get_num_devices)
 ialias (omp_is_initial_device)
 ialias (omp_get_device_num)
+ialias (omp_get_max_teams)
diff --git a/libgomp/config/nvptx/icv-device.c b/libgomp/config/nvptx/icv-device.c
index faf90f9..f744e75 100644
--- a/libgomp/config/nvptx/icv-device.c
+++ b/libgomp/config/nvptx/icv-device.c
@@ -28,6 +28,16 @@
 
 #include "libgomp.h"
 
+/* This is set to the ICV values of current GPU during device initialization,
+   when the offload image containing this libgomp portion is loaded.  */
+static volatile int GOMP_DEVICE_NUM_VAR;
+static volatile int GOMP_RUN_SCHED_VAR;
+static volatile int GOMP_RUN_SCHED_CHUNK_SIZE;
+static volatile int GOMP_DEFAULT_DEVICE_VAR;
+static volatile int GOMP_MAX_ACTIVE_LEVELS_VAR;
+static volatile omp_proc_bind_t GOMP_BIND_VAR;
+static volatile int GOMP_NTEAMS_VAR;
+
 void
 omp_set_default_device (int device_num __attribute__((unused)))
 {
@@ -36,7 +46,7 @@ omp_set_default_device (int device_num __attribute__((unused)))
 int
 omp_get_default_device (void)
 {
-  return 0;
+  return GOMP_DEFAULT_DEVICE_VAR;
 }
 
 int
@@ -58,19 +68,22 @@ omp_is_initial_device (void)
   return 0;
 }
 
-/* This is set to the device number of current GPU during device initialization,
-   when the offload image containing this libgomp portion is loaded.  */
-static volatile int GOMP_DEVICE_NUM_VAR;
-
 int
 omp_get_device_num (void)
 {
   return GOMP_DEVICE_NUM_VAR;
 }
 
+int
+omp_get_max_teams (void)
+{
+  return GOMP_NTEAMS_VAR;
+}
+
 ialias (omp_set_default_device)
 ialias (omp_get_default_device)
 ialias (omp_get_initial_device)
 ialias (omp_get_num_devices)
 ialias (omp_is_initial_device)
 ialias (omp_get_device_num)
+ialias (omp_get_max_teams)
diff --git a/libgomp/env.c b/libgomp/env.c
index 1c4ee89..74803ec 100644
--- a/libgomp/env.c
+++ b/libgomp/env.c
@@ -73,6 +73,74 @@ struct gomp_task_icv gomp_global_icv = {
   .target_data = NULL
 };
 
+/* The initial ICV values for the host.  */
+struct gomp_initial_icv_t gomp_initial_icv;
+
+/* Initial ICV values that were configured for the host and for all devices by
+   using environment variables like OMP_NUM_TEAMS_ALL.  */
+struct gomp_initial_icv_t gomp_initial_icv_all;
+
+/* Initial ICV values that were configured only for devices (not for the host)
+   by using environment variables like OMP_NUM_TEAMS_DEV.  */
+struct gomp_initial_icv_t gomp_initial_icv_dev;
+
+/* Returns the element of the list for the specified device number.  */
+struct gomp_icv_list*
+gomp_get_icv_list (struct gomp_icv_list **list, int device_num)
+{
+  struct gomp_icv_list *l = *list;
+  while (l != NULL)
+    {
+      if (l->device_num == device_num)
+	return l;
+      l = l->next;
+    }
+  return NULL;
+}
+
+void*
+gomp_get_icv_value_ptr (struct gomp_icv_list **list, int device_num)
+{
+  struct gomp_icv_list *l = gomp_get_icv_list (list, device_num);
+  if (l == NULL)
+    return NULL;
+  return l->value;
+}
+
+/* Lists for initial device-specific ICVs, i.e. ICVs that are configured for
+   particular devices (with environment variables like OMP_NUM_TEAMS_DEV_42). */
+struct gomp_icv_list *gomp_dyn_var_dev_list = NULL;
+struct gomp_icv_list *gomp_nthreads_var_dev_list = NULL;
+struct gomp_icv_list *gomp_nthreads_var_list_dev_list = NULL;
+struct gomp_icv_list *gomp_nthreads_var_list_len_dev_list = NULL;
+struct gomp_icv_list *gomp_run_sched_var_dev_list = NULL;
+struct gomp_icv_list *gomp_run_sched_chunk_size_dev_list = NULL;
+struct gomp_icv_list *gomp_nteams_var_dev_list = NULL;
+struct gomp_icv_list *gomp_thread_limit_var_dev_list = NULL;
+struct gomp_icv_list *gomp_max_active_levels_var_dev_list = NULL;
+struct gomp_icv_list *gomp_proc_bind_var_dev_list = NULL;
+struct gomp_icv_list *gomp_proc_bind_var_list_dev_list = NULL;
+struct gomp_icv_list *gomp_proc_bind_var_list_len_dev_list = NULL;
+struct gomp_icv_list *stacksize_dev_list = NULL;
+struct gomp_icv_list *wait_policy_dev_list = NULL;
+struct gomp_icv_list *teams_thread_limit_var_dev_list = NULL;
+
+/* Flags for non-global ICVs to store by which environment variables they are
+   affected.  */
+struct gomp_icv_flags_t gomp_initial_icv_flags = {
+  .nthreads_var = GOMP_ENV_VAR_SUFFIX_UNKNOWN,
+  .run_sched_var = GOMP_ENV_VAR_SUFFIX_UNKNOWN,
+  .run_sched_chunk_size = GOMP_ENV_VAR_SUFFIX_UNKNOWN,
+  .thread_limit_var = GOMP_ENV_VAR_SUFFIX_UNKNOWN,
+  .dyn_var = GOMP_ENV_VAR_SUFFIX_UNKNOWN,
+  .max_active_levels_var = GOMP_ENV_VAR_SUFFIX_UNKNOWN,
+  .bind_var = GOMP_ENV_VAR_SUFFIX_UNKNOWN,
+  .nteams_var = GOMP_ENV_VAR_SUFFIX_UNKNOWN,
+  .stacksize = GOMP_ENV_VAR_SUFFIX_UNKNOWN,
+  .wait_policy = GOMP_ENV_VAR_SUFFIX_UNKNOWN,
+  .teams_thread_limit_var = GOMP_ENV_VAR_SUFFIX_UNKNOWN
+};
+
 bool gomp_cancel_var = false;
 enum gomp_target_offload_t gomp_target_offload_var
   = GOMP_TARGET_OFFLOAD_DEFAULT;
@@ -106,16 +174,17 @@ static unsigned long stacksize = GOMP_DEFAULT_STACKSIZE;
 
 /* Parse the OMP_SCHEDULE environment variable.  */
 
-static void
-parse_schedule (void)
+static bool
+parse_schedule (const char *name, enum gomp_schedule_type *schedule,
+		int *chunk_size)
 {
   char *env, *end;
   unsigned long value;
   int monotonic = 0;
 
-  env = getenv ("OMP_SCHEDULE");
+  env = getenv (name);
   if (env == NULL)
-    return;
+    return false;
 
   while (isspace ((unsigned char) *env))
     ++env;
@@ -141,38 +210,37 @@ parse_schedule (void)
     }
   if (strncasecmp (env, "static", 6) == 0)
     {
-      gomp_global_icv.run_sched_var = GFS_STATIC;
+      *schedule = GFS_STATIC;
       env += 6;
     }
   else if (strncasecmp (env, "dynamic", 7) == 0)
     {
-      gomp_global_icv.run_sched_var = GFS_DYNAMIC;
+      *schedule = GFS_DYNAMIC;
       env += 7;
     }
   else if (strncasecmp (env, "guided", 6) == 0)
     {
-      gomp_global_icv.run_sched_var = GFS_GUIDED;
+      *schedule = GFS_GUIDED;
       env += 6;
     }
   else if (strncasecmp (env, "auto", 4) == 0)
     {
-      gomp_global_icv.run_sched_var = GFS_AUTO;
+      *schedule = GFS_AUTO;
       env += 4;
     }
   else
     goto unknown;
 
   if (monotonic == 1
-      || (monotonic == 0 && gomp_global_icv.run_sched_var == GFS_STATIC))
-    gomp_global_icv.run_sched_var |= GFS_MONOTONIC;
+      || (monotonic == 0 && *schedule == GFS_STATIC))
+    *schedule |= GFS_MONOTONIC;
 
   while (isspace ((unsigned char) *env))
     ++env;
   if (*env == '\0')
     {
-      gomp_global_icv.run_sched_chunk_size
-	= (gomp_global_icv.run_sched_var & ~GFS_MONOTONIC) != GFS_STATIC;
-      return;
+      *chunk_size = (*schedule & ~GFS_MONOTONIC) != GFS_STATIC;
+      return true;
     }
   if (*env++ != ',')
     goto unknown;
@@ -194,20 +262,19 @@ parse_schedule (void)
   if ((int)value != value)
     goto invalid;
 
-  if (value == 0
-      && (gomp_global_icv.run_sched_var & ~GFS_MONOTONIC) != GFS_STATIC)
+  if (value == 0 && (*schedule & ~GFS_MONOTONIC) != GFS_STATIC)
     value = 1;
-  gomp_global_icv.run_sched_chunk_size = value;
-  return;
+  *chunk_size = value;
+  return true;
 
  unknown:
   gomp_error ("Unknown value for environment variable OMP_SCHEDULE");
-  return;
+  return false;
 
  invalid:
   gomp_error ("Invalid value for chunk size in "
 	      "environment variable OMP_SCHEDULE");
-  return;
+  return false;
 }
 
 /* Parse an unsigned long environment variable.  Return true if one was
@@ -1034,15 +1101,18 @@ parse_boolean (const char *name, bool *value)
 
 /* Parse the OMP_WAIT_POLICY environment variable and return the value.  */
 
-static int
-parse_wait_policy (void)
+static bool
+parse_wait_policy (const char *name, int *pvalue)
 {
   const char *env;
   int ret = -1;
 
-  env = getenv ("OMP_WAIT_POLICY");
+  env = getenv (name);
   if (env == NULL)
-    return -1;
+  {
+    *pvalue = -1;
+    return false;
+  }
 
   while (isspace ((unsigned char) *env))
     ++env;
@@ -1061,9 +1131,13 @@ parse_wait_policy (void)
   while (isspace ((unsigned char) *env))
     ++env;
   if (*env == '\0')
-    return ret;
+    {
+      *pvalue = ret;
+      return true;
+    }
   gomp_error ("Invalid value for environment variable OMP_WAIT_POLICY");
-  return -1;
+  *pvalue = -1;
+  return false;
 }
 
 /* Parse the GOMP_CPU_AFFINITY environment varible.  Return true if one was
@@ -1251,62 +1325,56 @@ parse_gomp_openacc_dim (void)
     }
 }
 
-void
-omp_display_env (int verbose)
+/* Helper function for omp_display_env which prints the values of run_sched_var.
+   'device' can be 'host', 'dev', 'all' or a particular device number.  */
+static void
+print_schedule (enum gomp_schedule_type run_sched_var, int run_sched_chunk_size,
+		const char* device)
 {
-  int i;
-
-  fputs ("\nOPENMP DISPLAY ENVIRONMENT BEGIN\n", stderr);
-
-  fputs ("  _OPENMP = '201511'\n", stderr);
-  fprintf (stderr, "  OMP_DYNAMIC = '%s'\n",
-	   gomp_global_icv.dyn_var ? "TRUE" : "FALSE");
-  fprintf (stderr, "  OMP_NESTED = '%s'\n",
-	   gomp_global_icv.max_active_levels_var > 1 ? "TRUE" : "FALSE");
-
-  fprintf (stderr, "  OMP_NUM_THREADS = '%lu", gomp_global_icv.nthreads_var);
-  for (i = 1; i < gomp_nthreads_var_list_len; i++)
-    fprintf (stderr, ",%lu", gomp_nthreads_var_list[i]);
-  fputs ("'\n", stderr);
-
-  fprintf (stderr, "  OMP_SCHEDULE = '");
-  if ((gomp_global_icv.run_sched_var & GFS_MONOTONIC))
+  fprintf (stderr, "  [%s] OMP_SCHEDULE = '", device);
+  if ((run_sched_var & GFS_MONOTONIC))
     {
-      if (gomp_global_icv.run_sched_var != (GFS_MONOTONIC | GFS_STATIC))
+      if (run_sched_var != (GFS_MONOTONIC | GFS_STATIC))
 	fputs ("MONOTONIC:", stderr);
     }
-  else if (gomp_global_icv.run_sched_var == GFS_STATIC)
+  else if (run_sched_var == GFS_STATIC)
     fputs ("NONMONOTONIC:", stderr);
-  switch (gomp_global_icv.run_sched_var & ~GFS_MONOTONIC)
+  switch (run_sched_var & ~GFS_MONOTONIC)
     {
     case GFS_RUNTIME:
       fputs ("RUNTIME", stderr);
-      if (gomp_global_icv.run_sched_chunk_size != 1)
-	fprintf (stderr, ",%d", gomp_global_icv.run_sched_chunk_size);
+      if (run_sched_chunk_size != 1)
+	fprintf (stderr, ",%d", run_sched_chunk_size);
       break;
     case GFS_STATIC:
       fputs ("STATIC", stderr);
-      if (gomp_global_icv.run_sched_chunk_size != 0)
-	fprintf (stderr, ",%d", gomp_global_icv.run_sched_chunk_size);
+      if (run_sched_chunk_size != 0)
+	fprintf (stderr, ",%d", run_sched_chunk_size);
       break;
     case GFS_DYNAMIC:
       fputs ("DYNAMIC", stderr);
-      if (gomp_global_icv.run_sched_chunk_size != 1)
-	fprintf (stderr, ",%d", gomp_global_icv.run_sched_chunk_size);
+      if (run_sched_chunk_size != 1)
+	fprintf (stderr, ",%d", run_sched_chunk_size);
       break;
     case GFS_GUIDED:
       fputs ("GUIDED", stderr);
-      if (gomp_global_icv.run_sched_chunk_size != 1)
-	fprintf (stderr, ",%d", gomp_global_icv.run_sched_chunk_size);
+      if (run_sched_chunk_size != 1)
+	fprintf (stderr, ",%d", run_sched_chunk_size);
       break;
     case GFS_AUTO:
       fputs ("AUTO", stderr);
       break;
     }
   fputs ("'\n", stderr);
+}
 
-  fputs ("  OMP_PROC_BIND = '", stderr);
-  switch (gomp_global_icv.bind_var)
+/* Helper function for omp_display_env which prints the values of proc_bind_var.
+   'device' can be 'host', 'dev', 'all', or a particular device number.  */
+static void print_proc_bind (char proc_bind_var, unsigned long len, char **list,
+			     const char* device)
+{
+  fprintf (stderr, "  [%s] OMP_PROC_BIND = '", device);
+  switch (proc_bind_var)
     {
     case omp_proc_bind_false:
       fputs ("FALSE", stderr);
@@ -1324,8 +1392,8 @@ omp_display_env (int verbose)
       fputs ("SPREAD", stderr);
       break;
     }
-  for (i = 1; i < gomp_bind_var_list_len; i++)
-    switch (gomp_bind_var_list[i])
+  for (int i = 1; i < len; i++)
+    switch ((*list)[i])
       {
       case omp_proc_bind_master:
 	fputs (",MASTER", stderr); /* TODO: Change to PRIMARY for OpenMP 5.1. */
@@ -1338,7 +1406,131 @@ omp_display_env (int verbose)
 	break;
       }
   fputs ("'\n", stderr);
-  fputs ("  OMP_PLACES = '", stderr);
+}
+
+void
+omp_display_env (int verbose)
+{
+  int i;
+
+  fputs ("\nOPENMP DISPLAY ENVIRONMENT BEGIN\n", stderr);
+
+  fputs ("  _OPENMP = '201511'\n", stderr);
+
+  fprintf (stderr, "  [host] OMP_DYNAMIC = '%s'\n",
+	   gomp_initial_icv.dyn_var ? "TRUE" : "FALSE");
+  if (gomp_initial_icv_flags.dyn_var & GOMP_ENV_VAR_SUFFIX_ALL)
+    fprintf (stderr, "  [all] OMP_DYNAMIC = '%s'\n",
+	     gomp_initial_icv_all.dyn_var ? "TRUE" : "FALSE");
+  if (gomp_initial_icv_flags.dyn_var & GOMP_ENV_VAR_SUFFIX_DEV)
+    fprintf (stderr, "  [device] OMP_DYNAMIC = '%s'\n",
+	     gomp_initial_icv_dev.dyn_var ? "TRUE" : "FALSE");
+  struct gomp_icv_list* l_dyn_var = gomp_dyn_var_dev_list;
+  while (l_dyn_var != NULL)
+    {
+      fprintf (stderr, "  [%d] OMP_DYNAMIC = '%s'\n", l_dyn_var->device_num,
+	       *(bool*)l_dyn_var->value ? "TRUE" : "FALSE");
+      l_dyn_var = l_dyn_var->next;
+    }
+
+  /* The OMP_NESTED environment variable has been deprecated.  */
+  fprintf (stderr, "  [host] OMP_NESTED = '%s'\n",
+	   gomp_initial_icv.max_active_levels_var > 1 ? "TRUE" : "FALSE");
+
+  fprintf (stderr, "  [host] OMP_NUM_THREADS = '%lu",
+	   gomp_initial_icv.nthreads_var);
+  for (i = 1; i < gomp_nthreads_var_list_len; i++)
+    fprintf (stderr, ",%lu", gomp_nthreads_var_list[i]);
+  fputs ("'\n", stderr);
+  if (gomp_initial_icv_flags.nthreads_var & GOMP_ENV_VAR_SUFFIX_ALL)
+    {
+      fprintf (stderr, "  [all] OMP_NUM_THREADS = '%lu",
+	       gomp_initial_icv_all.nthreads_var);
+      for (i = 1; i < gomp_initial_icv_all.nthreads_var_list_len; i++)
+	fprintf (stderr, ",%lu", gomp_initial_icv_all.nthreads_var_list[i]);
+      fputs ("'\n", stderr);
+    }
+  if (gomp_initial_icv_flags.nthreads_var & GOMP_ENV_VAR_SUFFIX_DEV)
+    {
+      fprintf (stderr, "  [device] OMP_NUM_THREADS = '%lu",
+	       gomp_initial_icv_dev.nthreads_var);
+      for (i = 1; i < gomp_initial_icv_dev.nthreads_var_list_len; i++)
+	fprintf (stderr, ",%lu", gomp_initial_icv_dev.nthreads_var_list[i]);
+      fputs ("'\n", stderr);
+    }
+  struct gomp_icv_list* l_nthreads_var = gomp_nthreads_var_dev_list;
+  while (l_nthreads_var != NULL)
+    {
+      fprintf (stderr, "  [%d] OMP_NUM_THREADS = '%lu",
+	       l_nthreads_var->device_num,
+	       *(unsigned long*)l_nthreads_var->value);
+      struct gomp_icv_list *len
+	= gomp_get_icv_list (&gomp_nthreads_var_list_len_dev_list,
+			     l_nthreads_var->device_num);
+      if (len != NULL)
+	{
+	  struct gomp_icv_list *list
+	    = gomp_get_icv_list (&gomp_nthreads_var_list_dev_list,
+				 l_nthreads_var->device_num);
+	  for (i = 1; i < *(unsigned long*)len->value; i++)
+	    fprintf (stderr, ",%lu", (*(unsigned long**)list->value)[i]);
+	}
+      fputs ("'\n", stderr);
+      l_nthreads_var = l_nthreads_var->next;
+    }
+
+  print_schedule (gomp_initial_icv.run_sched_var,
+		  gomp_initial_icv.run_sched_chunk_size, "host");
+  if (gomp_initial_icv_flags.run_sched_var & GOMP_ENV_VAR_SUFFIX_ALL)
+    print_schedule (gomp_initial_icv_all.run_sched_var,
+		    gomp_initial_icv_all.run_sched_chunk_size, "all");
+  if (gomp_initial_icv_flags.run_sched_var & GOMP_ENV_VAR_SUFFIX_DEV)
+    print_schedule (gomp_initial_icv_dev.run_sched_var,
+		    gomp_initial_icv_dev.run_sched_chunk_size, "device");
+  struct gomp_icv_list* l_run_sched_var = gomp_run_sched_var_dev_list;
+  while (l_run_sched_var != NULL)
+    {
+      struct gomp_icv_list* l_run_sched_chunk_size
+	= gomp_get_icv_list (&gomp_run_sched_chunk_size_dev_list,
+			     l_run_sched_var->device_num);
+      int chunk_size = *(int*)l_run_sched_chunk_size->value;
+      char dev_num[10];
+      sprintf (dev_num, "%d", l_run_sched_var->device_num);
+      enum gomp_schedule_type schedule_type
+	= *(enum gomp_schedule_type*)l_run_sched_var->value;
+      print_schedule (schedule_type, chunk_size, dev_num);
+      l_run_sched_var = l_run_sched_var->next;
+    }
+
+  print_proc_bind (gomp_initial_icv.bind_var,
+		   gomp_initial_icv.bind_var_list_len,
+		   &gomp_initial_icv.bind_var_list, "host");
+  if (gomp_initial_icv_flags.bind_var & GOMP_ENV_VAR_SUFFIX_ALL)
+    print_proc_bind (gomp_initial_icv_all.bind_var,
+		     gomp_initial_icv_all.bind_var_list_len,
+		     &gomp_initial_icv_all.bind_var_list, "all");
+  if (gomp_initial_icv_flags.bind_var & GOMP_ENV_VAR_SUFFIX_DEV)
+    print_proc_bind (gomp_initial_icv_dev.bind_var,
+		     gomp_initial_icv_dev.bind_var_list_len,
+		     &gomp_initial_icv_dev.bind_var_list, "device");
+  struct gomp_icv_list* l_proc_bind_var = gomp_proc_bind_var_dev_list;
+  while (l_proc_bind_var != NULL)
+    {
+      struct gomp_icv_list *list
+	= gomp_get_icv_list (&gomp_proc_bind_var_list_dev_list,
+			     l_proc_bind_var->device_num);
+      struct gomp_icv_list *len
+	= gomp_get_icv_list (&gomp_proc_bind_var_list_len_dev_list,
+			     l_proc_bind_var->device_num);
+      char dev_num[10];
+      sprintf (dev_num, "%d", l_proc_bind_var->device_num);
+      char proc_bind = *(char*)l_proc_bind_var->value;
+      print_proc_bind (proc_bind, *(unsigned long*)len->value,
+		       &*(char**)list->value, dev_num);
+      l_proc_bind_var = l_proc_bind_var->next;
+    }
+
+  fputs ("  [host] OMP_PLACES = '", stderr);
   for (i = 0; i < gomp_places_list_len; i++)
     {
       fputs ("{", stderr);
@@ -1347,30 +1539,119 @@ omp_display_env (int verbose)
     }
   fputs ("'\n", stderr);
 
-  fprintf (stderr, "  OMP_STACKSIZE = '%lu'\n", stacksize);
+  fprintf (stderr, "  [host] OMP_STACKSIZE = '%lu'\n",
+	   gomp_initial_icv.stacksize);
+  if (gomp_initial_icv_flags.stacksize & GOMP_ENV_VAR_SUFFIX_ALL)
+    fprintf (stderr, "  [all] OMP_STACKSIZE = '%lu'\n",
+	     gomp_initial_icv_all.stacksize);
+  if (gomp_initial_icv_flags.stacksize & GOMP_ENV_VAR_SUFFIX_DEV)
+    fprintf (stderr, "  [device] OMP_STACKSIZE = '%lu'\n",
+	     gomp_initial_icv_dev.stacksize);
+  struct gomp_icv_list* l_stacksize = stacksize_dev_list;
+  while (l_stacksize != NULL)
+    {
+      fprintf (stderr, "  [%d] OMP_STACKSIZE = '%lu'\n",
+	       l_stacksize->device_num, *(unsigned long*)l_stacksize->value);
+      l_stacksize = l_stacksize->next;
+    }
 
   /* GOMP's default value is actually neither active nor passive.  */
-  fprintf (stderr, "  OMP_WAIT_POLICY = '%s'\n",
-	   wait_policy > 0 ? "ACTIVE" : "PASSIVE");
-  fprintf (stderr, "  OMP_THREAD_LIMIT = '%u'\n",
-	   gomp_global_icv.thread_limit_var);
-  fprintf (stderr, "  OMP_MAX_ACTIVE_LEVELS = '%u'\n",
-	   gomp_global_icv.max_active_levels_var);
-  fprintf (stderr, "  OMP_NUM_TEAMS = '%u'\n", gomp_nteams_var);
-  fprintf (stderr, "  OMP_TEAMS_THREAD_LIMIT = '%u'\n",
-	   gomp_teams_thread_limit_var);
-
-  fprintf (stderr, "  OMP_CANCELLATION = '%s'\n",
+  fprintf (stderr, "  [host] OMP_WAIT_POLICY = '%s'\n",
+	   gomp_initial_icv.wait_policy > 0 ? "ACTIVE" : "PASSIVE");
+  if (gomp_initial_icv_flags.wait_policy & GOMP_ENV_VAR_SUFFIX_ALL)
+    fprintf (stderr, "  [all] OMP_WAIT_POLICY = '%s'\n",
+	     gomp_initial_icv_all.wait_policy > 0 ? "ACTIVE" : "PASSIVE");
+  if (gomp_initial_icv_flags.wait_policy & GOMP_ENV_VAR_SUFFIX_DEV)
+    fprintf (stderr, "  [device] OMP_WAIT_POLICY = '%s'\n",
+	     gomp_initial_icv_dev.wait_policy > 0 ? "ACTIVE" : "PASSIVE");
+  struct gomp_icv_list* l_wait_policy = wait_policy_dev_list;
+  while (l_wait_policy != NULL)
+    {
+      fprintf (stderr, "  [%d] OMP_WAIT_POLICY = '%s'\n",
+	       l_wait_policy->device_num,
+	       *(int*)l_wait_policy->value > 0 ? "ACTIVE" : "PASSIVE");
+      l_wait_policy = l_wait_policy->next;
+    }
+
+  fprintf (stderr, "  [host] OMP_THREAD_LIMIT = '%u'\n",
+	   gomp_initial_icv.thread_limit_var);
+  if (gomp_initial_icv_flags.thread_limit_var & GOMP_ENV_VAR_SUFFIX_ALL)
+    fprintf (stderr, "  [all] OMP_THREAD_LIMIT = '%d'\n",
+	     gomp_initial_icv_all.thread_limit_var);
+  if (gomp_initial_icv_flags.thread_limit_var & GOMP_ENV_VAR_SUFFIX_DEV)
+    fprintf (stderr, "  [device] OMP_THREAD_LIMIT = '%d'\n",
+	     gomp_initial_icv_dev.thread_limit_var);
+  struct gomp_icv_list* l_thread_limit = gomp_thread_limit_var_dev_list;
+  while (l_thread_limit != NULL)
+    {
+      fprintf (stderr, "  [%d] OMP_THREAD_LIMIT = '%d'\n",
+	       l_thread_limit->device_num, *(int*)l_thread_limit->value);
+      l_thread_limit = l_thread_limit->next;
+    }
+
+  fprintf (stderr, "  [host] OMP_MAX_ACTIVE_LEVELS = '%u'\n",
+	   gomp_initial_icv.max_active_levels_var);
+  if (gomp_initial_icv_flags.max_active_levels_var & GOMP_ENV_VAR_SUFFIX_ALL)
+    fprintf (stderr, "  [all] OMP_MAX_ACTIVE_LEVELS = '%u'\n",
+	     gomp_initial_icv_all.max_active_levels_var);
+  if (gomp_initial_icv_flags.max_active_levels_var & GOMP_ENV_VAR_SUFFIX_DEV)
+    fprintf (stderr, "  [device] OMP_MAX_ACTIVE_LEVELS = '%u'\n",
+	     gomp_initial_icv_dev.max_active_levels_var);
+  struct gomp_icv_list* l_max_active_levels
+    = gomp_max_active_levels_var_dev_list;
+  while (l_max_active_levels != NULL)
+    {
+      fprintf (stderr, "  [%d] OMP_MAX_ACTIVE_LEVELS = '%u'\n",
+	       l_max_active_levels->device_num,
+	       *(int*)l_max_active_levels->value);
+      l_max_active_levels = l_max_active_levels->next;
+    }
+
+  fprintf (stderr, "  [host] OMP_NUM_TEAMS = '%d'\n",
+	   gomp_initial_icv.nteams_var);
+  if (gomp_initial_icv_flags.nteams_var & GOMP_ENV_VAR_SUFFIX_ALL)
+    fprintf (stderr, "  [all] OMP_NUM_TEAMS = '%d'\n",
+	     gomp_initial_icv_all.nteams_var);
+  if (gomp_initial_icv_flags.nteams_var & GOMP_ENV_VAR_SUFFIX_DEV)
+    fprintf (stderr, "  [device] OMP_NUM_TEAMS = '%d'\n",
+	     gomp_initial_icv_dev.nteams_var);
+  struct gomp_icv_list* l_nteams_var = gomp_nteams_var_dev_list;
+  while (l_nteams_var != NULL)
+    {
+      fprintf (stderr, "  [%d] OMP_NUM_TEAMS = '%d'\n",
+	       l_nteams_var->device_num, *(int*)l_nteams_var->value);
+      l_nteams_var = l_nteams_var->next;
+    }
+
+  fprintf (stderr, "  [host] OMP_TEAMS_THREAD_LIMIT = '%u'\n",
+	   gomp_initial_icv.teams_thread_limit_var);
+  if (gomp_initial_icv_flags.teams_thread_limit_var & GOMP_ENV_VAR_SUFFIX_ALL)
+    fprintf (stderr, "  [all] OMP_TEAMS_THREAD_LIMIT = '%u'\n",
+	     gomp_initial_icv_all.teams_thread_limit_var);
+  if (gomp_initial_icv_flags.teams_thread_limit_var & GOMP_ENV_VAR_SUFFIX_DEV)
+    fprintf (stderr, "  [device] OMP_TEAMS_THREAD_LIMIT = '%u'\n",
+	     gomp_initial_icv_dev.teams_thread_limit_var);
+  struct gomp_icv_list* l_teams_thr_limit = teams_thread_limit_var_dev_list;
+  while (l_teams_thr_limit != NULL)
+    {
+      fprintf (stderr, "  [%d] OMP_TEAMS_THREAD_LIMIT = '%u'\n",
+	       l_teams_thr_limit->device_num, *(int*)l_teams_thr_limit->value);
+      l_teams_thr_limit = l_teams_thr_limit->next;
+    }
+
+  fprintf (stderr, "  [all] OMP_CANCELLATION = '%s'\n",
 	   gomp_cancel_var ? "TRUE" : "FALSE");
-  fprintf (stderr, "  OMP_DEFAULT_DEVICE = '%d'\n",
-	   gomp_global_icv.default_device_var);
-  fprintf (stderr, "  OMP_MAX_TASK_PRIORITY = '%d'\n",
+
+  fprintf (stderr, "  [all] OMP_DEFAULT_DEVICE = '%d'\n",
+	   gomp_initial_icv.default_device_var);
+
+  fprintf (stderr, "  [all] OMP_MAX_TASK_PRIORITY = '%d'\n",
 	   gomp_max_task_priority_var);
-  fprintf (stderr, "  OMP_DISPLAY_AFFINITY = '%s'\n",
+  fprintf (stderr, "  [all] OMP_DISPLAY_AFFINITY = '%s'\n",
 	   gomp_display_affinity_var ? "TRUE" : "FALSE");
-  fprintf (stderr, "  OMP_AFFINITY_FORMAT = '%s'\n",
+  fprintf (stderr, "  [host] OMP_AFFINITY_FORMAT = '%s'\n",
 	   gomp_affinity_format_var);
-  fprintf (stderr, "  OMP_ALLOCATOR = '");
+  fprintf (stderr, "  [host] OMP_ALLOCATOR = '");
   switch (gomp_def_allocator)
     {
 #define C(v) case v: fputs (#v, stderr); break;
@@ -1387,7 +1668,7 @@ omp_display_env (int verbose)
     }
   fputs ("'\n", stderr);
 
-  fputs ("  OMP_TARGET_OFFLOAD = '", stderr);
+  fputs ("  [all] OMP_TARGET_OFFLOAD = '", stderr);
   switch (gomp_target_offload_var)
     {
     case GOMP_TARGET_OFFLOAD_DEFAULT:
@@ -1404,13 +1685,13 @@ omp_display_env (int verbose)
 
   if (verbose)
     {
-      fputs ("  GOMP_CPU_AFFINITY = ''\n", stderr);
-      fprintf (stderr, "  GOMP_STACKSIZE = '%lu'\n", stacksize);
+      fputs ("  [host] GOMP_CPU_AFFINITY = ''\n", stderr);
+      fprintf (stderr, "  [host] GOMP_STACKSIZE = '%lu'\n", stacksize);
 #ifdef HAVE_INTTYPES_H
-      fprintf (stderr, "  GOMP_SPINCOUNT = '%"PRIu64"'\n",
+      fprintf (stderr, "  [host] GOMP_SPINCOUNT = '%"PRIu64"'\n",
 	       (uint64_t) gomp_spin_count_var);
 #else
-      fprintf (stderr, "  GOMP_SPINCOUNT = '%lu'\n",
+      fprintf (stderr, "  [host] GOMP_SPINCOUNT = '%lu'\n",
 	       (unsigned long) gomp_spin_count_var);
 #endif
     }
@@ -1459,6 +1740,249 @@ handle_omp_display_env (void)
     ialias_call (omp_display_env) (verbose);
 }
 
+/* Helper function for parse_device_specific. Extracts the device number from
+   an environment variable name.  */
+static void
+get_device_num (char **env, int prefix_len, int *dev_num, int *name_len)
+{
+  if (env == NULL || *env == NULL)
+    {
+      *name_len = 0;
+      *dev_num = -1;
+      return;
+    }
+
+  int eq_pos = strchr (*env, '=') - *env;
+  int dev_num_len = eq_pos - prefix_len;
+  char buf_dev_num[dev_num_len+1];
+
+  strncpy(buf_dev_num, *env + prefix_len, dev_num_len);
+  buf_dev_num[dev_num_len] = '\0';
+  *dev_num = atoi (buf_dev_num);
+  *name_len = eq_pos;
+}
+
+/* Helper function for parse_device_specific. Adds a new node to the given
+   list.  */
+static struct gomp_icv_list*
+add_device_specific_icv (int dev_num, size_t size, struct gomp_icv_list **list)
+{
+  if (list == NULL)
+    return NULL;
+
+  struct gomp_icv_list *new_node =
+    (struct gomp_icv_list*) malloc (sizeof (struct gomp_icv_list));
+  new_node->device_num = dev_num;
+  new_node->value = malloc (size);
+  new_node->next = *list;
+  *list = new_node;
+
+  return new_node;
+}
+
+/* Helper function for 'initialize_env' to parse device-specific environment
+   variables like 'OMP_NUM_TEAMS_DEV_42'.  */
+static void
+parse_device_specific ()
+{
+  extern char **environ;
+  int dev_num;
+  int name_len;
+  struct gomp_icv_list *new_node;
+
+  for (char **env = environ; *env != 0; env++)
+  {
+    if (strncmp (*env, "OMP_SCHEDULE_DEV_", 17) == 0)
+      {
+	get_device_num (env, 17, &dev_num, &name_len);
+	char name[name_len+1];
+	strncpy(name, *env, name_len);
+	name[name_len] = '\0';
+	enum gomp_schedule_type schedule_type;
+	int chunk_size;
+	parse_schedule (name, &schedule_type, &chunk_size);
+	new_node = add_device_specific_icv (dev_num,
+					    sizeof (enum gomp_schedule_type),
+					    &gomp_run_sched_var_dev_list);
+	*((enum gomp_schedule_type*)(new_node->value)) = schedule_type;
+	new_node = add_device_specific_icv (dev_num, sizeof (int),
+					    &gomp_run_sched_chunk_size_dev_list);
+	*((int*)(new_node->value)) = chunk_size;
+	goto next;
+      }
+    else if (strncmp (*env, "OMP_DYNAMIC_DEV_", 16) == 0)
+      {
+	get_device_num (env, 16, &dev_num, &name_len);
+	char name[name_len+1];
+	strncpy(name, *env, name_len);
+	name[name_len] = '\0';
+	bool value;
+	if (!parse_boolean (name, &value))
+	  continue;
+	new_node = add_device_specific_icv (dev_num, sizeof (bool),
+					    &gomp_dyn_var_dev_list);
+	*((bool*)(new_node->value)) = value;
+	goto next;
+      }
+    else if (strncmp (*env, "OMP_THREAD_LIMIT_DEV_", 21) == 0)
+      {
+	get_device_num (env, 21, &dev_num, &name_len);
+	char name[name_len+1];
+	strncpy(name, *env, name_len);
+	name[name_len] = '\0';
+	unsigned long value;
+	if (!parse_unsigned_long (name, &value, false))
+	  continue;
+	value = value > INT_MAX ? UINT_MAX : value;
+	new_node = add_device_specific_icv (dev_num, sizeof (unsigned long),
+					    &gomp_thread_limit_var_dev_list);
+	*((unsigned long*)(new_node->value)) = value;
+	goto next;
+      }
+    else if (strncmp (*env, "OMP_NUM_THREADS_DEV_", 20) == 0)
+      {
+	get_device_num (env, 20, &dev_num, &name_len);
+	char name[name_len+1];
+	strncpy(name, *env, name_len);
+	name[name_len] = '\0';
+	unsigned long value;
+	unsigned long *pvalues;
+	unsigned long nvalues = 0;
+	if (!parse_unsigned_long_list (name, &value, &pvalues, &nvalues))
+	  continue;
+	new_node = add_device_specific_icv (dev_num, sizeof (unsigned long),
+					    &gomp_nthreads_var_dev_list);
+	*((unsigned long*)(new_node->value)) = value;
+	if (nvalues > 0)
+	  {
+	    new_node = add_device_specific_icv (dev_num, sizeof (unsigned long*),
+						&gomp_nthreads_var_list_dev_list);
+	    *((unsigned long**)(new_node->value)) = pvalues;
+
+	    new_node = add_device_specific_icv (dev_num, sizeof (unsigned long*),
+						&gomp_nthreads_var_list_len_dev_list);
+	    *((unsigned long*)(new_node->value)) = nvalues;
+	  }
+	goto next;
+      }
+    else if (strncmp (*env, "OMP_NUM_TEAMS_DEV_", 18) == 0)
+      {
+	get_device_num (env, 18, &dev_num, &name_len);
+	char name[name_len+1];
+	strncpy(name, *env, name_len);
+	name[name_len] = '\0';
+	int value;
+	if (!parse_int (name, &value, false))
+	  continue;
+	new_node = add_device_specific_icv (dev_num, sizeof (int),
+					    &gomp_nteams_var_dev_list);
+	*((int*)(new_node->value)) = value;
+	goto next;
+      }
+    else if (strncmp (*env, "OMP_PROC_BIND_DEV_", 18) == 0)
+      {
+	get_device_num (env, 18, &dev_num, &name_len);
+	char name[name_len+1];
+	strncpy(name, *env, name_len);
+	name[name_len] = '\0';
+	char value = omp_proc_bind_false;
+	char *pvalues;
+	unsigned long nvalues = 0;
+	bool ignore = false;
+	if (parse_bind_var (name, &value, &pvalues, &nvalues))
+	  ignore = true;
+
+	char omp_places[name_len];
+	strncpy(omp_places, *env, name_len);
+	if (parse_places_var (omp_places, ignore))
+	  {
+	    if (value == omp_proc_bind_false)
+	      value = true;
+	    ignore = true;
+	  }
+	if (parse_affinity (ignore))
+	  if (value == omp_proc_bind_false)
+	    value = true;
+
+	new_node = add_device_specific_icv (dev_num, sizeof (char),
+					    &gomp_proc_bind_var_dev_list);
+	*((char*)(new_node->value)) = value;
+
+	if (nvalues > 0)
+	  {
+	    new_node = add_device_specific_icv (dev_num, sizeof (char*),
+						&gomp_proc_bind_var_list_dev_list);
+	    *((char**)(new_node->value)) = pvalues;
+
+	    new_node = add_device_specific_icv (dev_num, sizeof (unsigned long),
+						&gomp_proc_bind_var_list_len_dev_list);
+	    *((unsigned long*)(new_node->value)) = nvalues;
+	  }
+	goto next;
+      }
+    else if (strncmp (*env, "OMP_MAX_ACTIVE_LEVELS_DEV_", 26) == 0)
+      {
+	get_device_num (env, 26, &dev_num, &name_len);
+	char name[name_len+1];
+	strncpy(name, *env, name_len);
+	name[name_len] = '\0';
+	unsigned long value;
+	if (!parse_unsigned_long (name, &value, true))
+	  continue;
+	value = (value > gomp_supported_active_levels)
+		? gomp_supported_active_levels : value;
+	new_node = add_device_specific_icv (dev_num, sizeof (unsigned long),
+					    &gomp_max_active_levels_var_dev_list);
+	*((unsigned long*)(new_node->value)) = value;
+	goto next;
+      }
+    else if (strncmp (*env, "OMP_STACKSIZE_DEV_", 18) == 0)
+      {
+	get_device_num (env, 18, &dev_num, &name_len);
+	char name[name_len+1];
+	strncpy(name, *env, name_len);
+	name[name_len] = '\0';
+	unsigned long value;
+	if (!parse_stacksize (name, &value))
+	  continue;
+	new_node = add_device_specific_icv (dev_num, sizeof (unsigned long),
+					    &stacksize_dev_list);
+	*((unsigned long*)(new_node->value)) = value;
+	goto next;
+      }
+    else if (strncmp (*env, "OMP_WAIT_POLICY_DEV_", 20) == 0)
+      {
+	get_device_num (env, 20, &dev_num, &name_len);
+	char name[name_len+1];
+	strncpy(name, *env, name_len);
+	name[name_len] = '\0';
+	int value;
+	if (!parse_wait_policy (name, &value))
+	  continue;
+	new_node = add_device_specific_icv (dev_num, sizeof (int),
+					    &wait_policy_dev_list);
+	*((int*)(new_node->value)) = value;
+	goto next;
+      }
+    else if (strncmp (*env, "OMP_TEAMS_THREAD_LIMIT_DEV_", 27) == 0)
+      {
+	get_device_num (env, 27, &dev_num, &name_len);
+	char name[name_len+1];
+	strncpy(name, *env, name_len);
+	name[name_len] = '\0';
+	int value;
+	if (!parse_int (name, &value, false))
+	  continue;
+	new_node = add_device_specific_icv (dev_num, sizeof (int),
+					    &teams_thread_limit_var_dev_list);
+	*((int*)(new_node->value)) = value;
+	goto next;
+      }
+
+ next:
+    new_node = NULL;
+  }
+}
 
 static void __attribute__((constructor))
 initialize_env (void)
@@ -1469,44 +1993,212 @@ initialize_env (void)
   /* Do a compile time check that mkomp_h.pl did good job.  */
   omp_check_defines ();
 
-  parse_schedule ();
-  parse_boolean ("OMP_DYNAMIC", &gomp_global_icv.dyn_var);
+  parse_device_specific ();
+
+  /* Parse the environment variables and store their values in the initial
+     struct.  */
+  if (parse_schedule ("OMP_SCHEDULE", &gomp_initial_icv.run_sched_var,
+		      &gomp_initial_icv.run_sched_chunk_size))
+    {
+      gomp_initial_icv_flags.run_sched_var |= GOMP_ENV_VAR_SUFFIX_NONE;
+      gomp_initial_icv_flags.run_sched_chunk_size |= GOMP_ENV_VAR_SUFFIX_NONE;
+    }
+  if (parse_schedule ("OMP_SCHEDULE_ALL", &gomp_initial_icv_all.run_sched_var,
+		      &gomp_initial_icv_all.run_sched_chunk_size))
+    {
+      gomp_initial_icv_flags.run_sched_var |= GOMP_ENV_VAR_SUFFIX_ALL;
+      gomp_initial_icv_flags.run_sched_chunk_size |= GOMP_ENV_VAR_SUFFIX_ALL;
+    }
+  if (parse_schedule ("OMP_SCHEDULE_DEV", &gomp_initial_icv_dev.run_sched_var,
+		      &gomp_initial_icv_dev.run_sched_chunk_size))
+    {
+      gomp_initial_icv_flags.run_sched_var |= GOMP_ENV_VAR_SUFFIX_DEV;
+      gomp_initial_icv_flags.run_sched_chunk_size |= GOMP_ENV_VAR_SUFFIX_DEV;
+    }
+
+  /* Set the ICV values for the host.  */
+  if (gomp_initial_icv_flags.run_sched_var & GOMP_ENV_VAR_SUFFIX_NONE)
+    {
+      gomp_global_icv.run_sched_var = gomp_initial_icv.run_sched_var;
+      gomp_global_icv.run_sched_chunk_size =
+	gomp_initial_icv.run_sched_chunk_size;
+    }
+  else if (gomp_initial_icv_flags.run_sched_var & GOMP_ENV_VAR_SUFFIX_ALL)
+    {
+      gomp_global_icv.run_sched_var = gomp_initial_icv_all.run_sched_var;
+      gomp_global_icv.run_sched_chunk_size =
+	gomp_initial_icv_all.run_sched_chunk_size;
+    }
+
+  if (parse_boolean ("OMP_DYNAMIC", &gomp_initial_icv.dyn_var))
+    gomp_initial_icv_flags.dyn_var |= GOMP_ENV_VAR_SUFFIX_NONE;
+  if (parse_boolean ("OMP_DYNAMIC_ALL", &gomp_initial_icv_all.dyn_var))
+    gomp_initial_icv_flags.dyn_var |= GOMP_ENV_VAR_SUFFIX_ALL;
+  if (parse_boolean ("OMP_DYNAMIC_DEV", &gomp_initial_icv_dev.dyn_var))
+    gomp_initial_icv_flags.dyn_var |= GOMP_ENV_VAR_SUFFIX_DEV;
+  if (gomp_initial_icv_flags.dyn_var & GOMP_ENV_VAR_SUFFIX_NONE)
+    gomp_global_icv.dyn_var = gomp_initial_icv.dyn_var;
+  else if (gomp_initial_icv_flags.dyn_var & GOMP_ENV_VAR_SUFFIX_ALL)
+    gomp_global_icv.dyn_var = gomp_initial_icv_all.dyn_var;
+
   parse_boolean ("OMP_CANCELLATION", &gomp_cancel_var);
   parse_boolean ("OMP_DISPLAY_AFFINITY", &gomp_display_affinity_var);
-  parse_int ("OMP_DEFAULT_DEVICE", &gomp_global_icv.default_device_var, true);
+
+  if (parse_int ("OMP_DEFAULT_DEVICE", &gomp_initial_icv.default_device_var,
+		 true))
+    gomp_global_icv.default_device_var = gomp_initial_icv.default_device_var;
+
   parse_target_offload ("OMP_TARGET_OFFLOAD", &gomp_target_offload_var);
   parse_int ("OMP_MAX_TASK_PRIORITY", &gomp_max_task_priority_var, true);
   gomp_def_allocator = parse_allocator ();
   if (parse_unsigned_long ("OMP_THREAD_LIMIT", &thread_limit_var, false))
     {
-      gomp_global_icv.thread_limit_var
+      gomp_initial_icv.thread_limit_var
+	= thread_limit_var > INT_MAX ? UINT_MAX : thread_limit_var;
+      gomp_initial_icv_flags.thread_limit_var |= GOMP_ENV_VAR_SUFFIX_NONE;
+    }
+  if (parse_unsigned_long ("OMP_THREAD_LIMIT_ALL", &thread_limit_var, false))
+    {
+      gomp_initial_icv_all.thread_limit_var
 	= thread_limit_var > INT_MAX ? UINT_MAX : thread_limit_var;
+      gomp_initial_icv_flags.thread_limit_var |= GOMP_ENV_VAR_SUFFIX_ALL;
     }
+  if (parse_unsigned_long ("OMP_THREAD_LIMIT_DEV", &thread_limit_var, false))
+    {
+      gomp_initial_icv_dev.thread_limit_var
+	= thread_limit_var > INT_MAX ? UINT_MAX : thread_limit_var;
+      gomp_initial_icv_flags.thread_limit_var |= GOMP_ENV_VAR_SUFFIX_DEV;
+    }
+  if (gomp_initial_icv_flags.thread_limit_var & GOMP_ENV_VAR_SUFFIX_NONE)
+    gomp_global_icv.thread_limit_var = gomp_initial_icv.thread_limit_var;
+  else if (gomp_initial_icv_flags.thread_limit_var & GOMP_ENV_VAR_SUFFIX_ALL)
+    gomp_global_icv.thread_limit_var = gomp_initial_icv_all.thread_limit_var;
+
   parse_int_secure ("GOMP_DEBUG", &gomp_debug_var, true);
 #ifndef HAVE_SYNC_BUILTINS
   gomp_mutex_init (&gomp_managed_threads_lock);
 #endif
   gomp_init_num_threads ();
   gomp_available_cpus = gomp_global_icv.nthreads_var;
-  if (!parse_unsigned_long_list ("OMP_NUM_THREADS",
-				 &gomp_global_icv.nthreads_var,
-				 &gomp_nthreads_var_list,
-				 &gomp_nthreads_var_list_len))
+
+  if (parse_unsigned_long_list ("OMP_NUM_THREADS",
+				&gomp_initial_icv.nthreads_var,
+				&gomp_initial_icv.nthreads_var_list,
+				&gomp_initial_icv.nthreads_var_list_len))
+    gomp_initial_icv_flags.nthreads_var |= GOMP_ENV_VAR_SUFFIX_NONE;
+  if (parse_unsigned_long_list ("OMP_NUM_THREADS_ALL",
+				&gomp_initial_icv_all.nthreads_var,
+				&gomp_initial_icv_all.nthreads_var_list,
+				&gomp_initial_icv_all.nthreads_var_list_len))
+    gomp_initial_icv_flags.nthreads_var |= GOMP_ENV_VAR_SUFFIX_ALL;
+  if (parse_unsigned_long_list ("OMP_NUM_THREADS_DEV",
+				&gomp_initial_icv_dev.nthreads_var,
+				&gomp_initial_icv_dev.nthreads_var_list,
+				&gomp_initial_icv_dev.nthreads_var_list_len))
+    gomp_initial_icv_flags.nthreads_var |= GOMP_ENV_VAR_SUFFIX_DEV;
+  if (gomp_initial_icv_flags.nthreads_var & GOMP_ENV_VAR_SUFFIX_NONE)
+    {
+      gomp_global_icv.nthreads_var = gomp_initial_icv.nthreads_var;
+      gomp_nthreads_var_list = gomp_initial_icv.nthreads_var_list;
+      gomp_nthreads_var_list_len = gomp_initial_icv.nthreads_var_list_len;
+    }
+  else if (gomp_initial_icv_flags.nthreads_var & GOMP_ENV_VAR_SUFFIX_ALL)
+    {
+      gomp_global_icv.nthreads_var = gomp_initial_icv_all.nthreads_var;
+      gomp_nthreads_var_list = gomp_initial_icv_all.nthreads_var_list;
+      gomp_nthreads_var_list_len = gomp_initial_icv_all.nthreads_var_list_len;
+    }
+  else
     gomp_global_icv.nthreads_var = gomp_available_cpus;
-  parse_int ("OMP_NUM_TEAMS", &gomp_nteams_var, false);
-  parse_int ("OMP_TEAMS_THREAD_LIMIT", &gomp_teams_thread_limit_var, false);
+
+  if (parse_int ("OMP_NUM_TEAMS", &gomp_initial_icv.nteams_var, false))
+    gomp_initial_icv_flags.nteams_var |= GOMP_ENV_VAR_SUFFIX_NONE;
+  if (parse_int ("OMP_NUM_TEAMS_ALL", &gomp_initial_icv_all.nteams_var, false))
+    gomp_initial_icv_flags.nteams_var |= GOMP_ENV_VAR_SUFFIX_ALL;
+  if (parse_int ("OMP_NUM_TEAMS_DEV", &gomp_initial_icv_dev.nteams_var, false))
+    gomp_initial_icv_flags.nteams_var |= GOMP_ENV_VAR_SUFFIX_DEV;
+  if (gomp_initial_icv_flags.nteams_var & GOMP_ENV_VAR_SUFFIX_NONE)
+    gomp_nteams_var = gomp_initial_icv.nteams_var;
+  else if (gomp_initial_icv_flags.nteams_var & GOMP_ENV_VAR_SUFFIX_ALL)
+    gomp_nteams_var = gomp_initial_icv_all.nteams_var;
+
+  if (parse_int ("OMP_TEAMS_THREAD_LIMIT",
+		 &gomp_initial_icv.teams_thread_limit_var, false))
+    gomp_initial_icv_flags.teams_thread_limit_var |= GOMP_ENV_VAR_SUFFIX_NONE;
+  if (parse_int ("OMP_TEAMS_THREAD_LIMIT_ALL",
+		 &gomp_initial_icv_all.teams_thread_limit_var, false))
+    gomp_initial_icv_flags.teams_thread_limit_var |= GOMP_ENV_VAR_SUFFIX_ALL;
+  if (parse_int ("OMP_TEAMS_THREAD_LIMIT_DEV",
+		 &gomp_initial_icv_dev.teams_thread_limit_var, false))
+    gomp_initial_icv_flags.teams_thread_limit_var |= GOMP_ENV_VAR_SUFFIX_DEV;
+  if (gomp_initial_icv_flags.teams_thread_limit_var & GOMP_ENV_VAR_SUFFIX_NONE)
+    gomp_teams_thread_limit_var = gomp_initial_icv.teams_thread_limit_var;
+  else if (gomp_initial_icv_flags.teams_thread_limit_var
+	   & GOMP_ENV_VAR_SUFFIX_ALL)
+    gomp_teams_thread_limit_var = gomp_initial_icv_all.teams_thread_limit_var;
+
   bool ignore = false;
-  if (parse_bind_var ("OMP_PROC_BIND",
-		      &gomp_global_icv.bind_var,
-		      &gomp_bind_var_list,
-		      &gomp_bind_var_list_len)
-      && gomp_global_icv.bind_var == omp_proc_bind_false)
-    ignore = true;
+  if (parse_bind_var ("OMP_PROC_BIND", &gomp_initial_icv.bind_var,
+		      &gomp_initial_icv.bind_var_list,
+		      &gomp_initial_icv.bind_var_list_len))
+    gomp_initial_icv_flags.bind_var |= GOMP_ENV_VAR_SUFFIX_NONE;
+  if (parse_bind_var ("OMP_PROC_BIND_ALL", &gomp_initial_icv_all.bind_var,
+		      &gomp_initial_icv_all.bind_var_list,
+		      &gomp_initial_icv_all.bind_var_list_len))
+    gomp_initial_icv_flags.bind_var |= GOMP_ENV_VAR_SUFFIX_ALL;
+  if (parse_bind_var ("OMP_PROC_BIND_DEV",
+		      &gomp_initial_icv_dev.bind_var,
+		      &gomp_initial_icv_dev.bind_var_list,
+		      &gomp_initial_icv_dev.bind_var_list_len))
+    gomp_initial_icv_flags.bind_var |= GOMP_ENV_VAR_SUFFIX_DEV;
+  if (gomp_initial_icv_flags.bind_var & GOMP_ENV_VAR_SUFFIX_NONE)
+    {
+      gomp_global_icv.bind_var = gomp_initial_icv.bind_var;
+      gomp_bind_var_list = gomp_initial_icv.bind_var_list;
+      gomp_bind_var_list_len = gomp_initial_icv.bind_var_list_len;
+      if (gomp_global_icv.bind_var == omp_proc_bind_false)
+	ignore = true;
+    }
+  else if (gomp_initial_icv_flags.bind_var & GOMP_ENV_VAR_SUFFIX_ALL)
+    {
+      gomp_global_icv.bind_var = gomp_initial_icv_all.bind_var;
+      gomp_bind_var_list = gomp_initial_icv_all.bind_var_list;
+      gomp_bind_var_list_len = gomp_initial_icv_all.bind_var_list_len;
+      if (gomp_global_icv.bind_var == omp_proc_bind_false)
+	ignore = true;
+    }
+
   if (parse_unsigned_long ("OMP_MAX_ACTIVE_LEVELS",
 			   &max_active_levels_var, true))
+    {
+      gomp_initial_icv.max_active_levels_var
+	= (max_active_levels_var > gomp_supported_active_levels)
+	  ? gomp_supported_active_levels : max_active_levels_var;
+      gomp_initial_icv_flags.max_active_levels_var |= GOMP_ENV_VAR_SUFFIX_NONE;
+    }
+  if (parse_unsigned_long ("OMP_MAX_ACTIVE_LEVELS_ALL",
+			   &max_active_levels_var, true))
+    {
+      gomp_initial_icv_all.max_active_levels_var
+	= (max_active_levels_var > gomp_supported_active_levels)
+	  ? gomp_supported_active_levels : max_active_levels_var;
+      gomp_initial_icv_flags.max_active_levels_var |= GOMP_ENV_VAR_SUFFIX_ALL;
+    }
+  if (parse_unsigned_long ("OMP_MAX_ACTIVE_LEVELS_DEV",
+			   &max_active_levels_var, true))
+    {
+      gomp_initial_icv_dev.max_active_levels_var
+	= (max_active_levels_var > gomp_supported_active_levels)
+	  ? gomp_supported_active_levels : max_active_levels_var;
+      gomp_initial_icv_flags.max_active_levels_var |= GOMP_ENV_VAR_SUFFIX_DEV;
+    }
+  if (gomp_initial_icv_flags.max_active_levels_var & GOMP_ENV_VAR_SUFFIX_NONE)
     gomp_global_icv.max_active_levels_var
-      = (max_active_levels_var > gomp_supported_active_levels)
-	? gomp_supported_active_levels : max_active_levels_var;
+      = gomp_initial_icv.max_active_levels_var;
+  else if (gomp_initial_icv_flags.max_active_levels_var
+	   & GOMP_ENV_VAR_SUFFIX_ALL)
+    gomp_global_icv.max_active_levels_var
+      = gomp_initial_icv_all.max_active_levels_var;
   else
     {
       bool nested = true;
@@ -1547,7 +2239,19 @@ initialize_env (void)
       gomp_set_affinity_format (env, strlen (env));
   }
 
-  wait_policy = parse_wait_policy ();
+  if (parse_wait_policy ("OMP_WAIT_POLICY", &gomp_initial_icv.wait_policy))
+    gomp_initial_icv_flags.wait_policy |= GOMP_ENV_VAR_SUFFIX_NONE;
+  if (parse_wait_policy ("OMP_WAIT_POLICY_ALL",
+			 &gomp_initial_icv_all.wait_policy))
+    gomp_initial_icv_flags.wait_policy |= GOMP_ENV_VAR_SUFFIX_ALL;
+  if (parse_wait_policy ("OMP_WAIT_POLICY_DEV",
+			 &gomp_initial_icv_dev.wait_policy))
+    gomp_initial_icv_flags.wait_policy |= GOMP_ENV_VAR_SUFFIX_DEV;
+  if (gomp_initial_icv_flags.wait_policy & GOMP_ENV_VAR_SUFFIX_NONE)
+    wait_policy = gomp_initial_icv.wait_policy;
+  else if (gomp_initial_icv_flags.wait_policy & GOMP_ENV_VAR_SUFFIX_ALL)
+    wait_policy = gomp_initial_icv_all.wait_policy;
+
   if (!parse_spincount ("GOMP_SPINCOUNT", &gomp_spin_count_var))
     {
       /* Using a rough estimation of 100000 spins per msec,
@@ -1573,8 +2277,20 @@ initialize_env (void)
   /* Not strictly environment related, but ordering constructors is tricky.  */
   pthread_attr_init (&gomp_thread_attr);
 
-  if (parse_stacksize ("OMP_STACKSIZE", &stacksize)
-      || parse_stacksize ("GOMP_STACKSIZE", &stacksize)
+  if (parse_stacksize ("OMP_STACKSIZE", &gomp_initial_icv.stacksize)
+      || parse_stacksize ("GOMP_STACKSIZE", &gomp_initial_icv.stacksize))
+    gomp_initial_icv_flags.stacksize |= GOMP_ENV_VAR_SUFFIX_NONE;
+  if (parse_stacksize ("OMP_STACKSIZE_ALL", &gomp_initial_icv_all.stacksize))
+    gomp_initial_icv_flags.stacksize |= GOMP_ENV_VAR_SUFFIX_ALL;
+  if (parse_stacksize ("OMP_STACKSIZE_DEV", &gomp_initial_icv_dev.stacksize))
+    gomp_initial_icv_flags.stacksize |= GOMP_ENV_VAR_SUFFIX_DEV;
+  if (gomp_initial_icv_flags.stacksize & GOMP_ENV_VAR_SUFFIX_NONE)
+    stacksize = gomp_initial_icv.stacksize;
+  else if (gomp_initial_icv_flags.stacksize & GOMP_ENV_VAR_SUFFIX_ALL)
+    stacksize = gomp_initial_icv_all.stacksize;
+
+  if (gomp_initial_icv_flags.stacksize & GOMP_ENV_VAR_SUFFIX_NONE
+      || gomp_initial_icv_flags.stacksize & GOMP_ENV_VAR_SUFFIX_ALL
       || GOMP_DEFAULT_STACKSIZE)
     {
       int err;
diff --git a/libgomp/icv-device.c b/libgomp/icv-device.c
index 79261ab..bc1af97 100644
--- a/libgomp/icv-device.c
+++ b/libgomp/icv-device.c
@@ -80,3 +80,11 @@ omp_get_device_num (void)
 }
 
 ialias (omp_get_device_num)
+
+int
+omp_get_max_teams (void)
+{
+  return gomp_nteams_var;
+}
+
+ialias (omp_get_max_teams)
diff --git a/libgomp/icv.c b/libgomp/icv.c
index de15cc8..5c435d6 100644
--- a/libgomp/icv.c
+++ b/libgomp/icv.c
@@ -155,12 +155,6 @@ omp_set_num_teams (int num_teams)
     gomp_nteams_var = num_teams;
 }
 
-int
-omp_get_max_teams (void)
-{
-  return gomp_nteams_var;
-}
-
 void
 omp_set_teams_thread_limit (int thread_limit)
 {
@@ -275,7 +269,6 @@ ialias (omp_set_max_active_levels)
 ialias (omp_get_max_active_levels)
 ialias (omp_get_supported_active_levels)
 ialias (omp_set_num_teams)
-ialias (omp_get_max_teams)
 ialias (omp_set_teams_thread_limit)
 ialias (omp_get_teams_thread_limit)
 ialias (omp_get_cancellation)
diff --git a/libgomp/libgomp-plugin.h b/libgomp/libgomp-plugin.h
index 07ab700..072cc47 100644
--- a/libgomp/libgomp-plugin.h
+++ b/libgomp/libgomp-plugin.h
@@ -102,11 +102,20 @@ struct addr_pair
   uintptr_t end;
 };
 
-/* This symbol is to name a target side variable that holds the designated
-   'device number' of the target device. The symbol needs to be available to
+/* This following symbols are to name target side variables that hold the
+   designated ICVs of the target device. The symbols need to be available to
    libgomp code and the offload plugin (which in the latter case must be
    stringified).  */
 #define GOMP_DEVICE_NUM_VAR __gomp_device_num
+#define GOMP_NTHREADS_VAR __gomp_nthreads
+#define GOMP_THREAD_LIMIT_VAR __gomp_thread_limit
+#define GOMP_RUN_SCHED_VAR __gomp_run_sched
+#define GOMP_RUN_SCHED_CHUNK_SIZE __gomp_run_sched_chunk_size
+#define GOMP_DEFAULT_DEVICE_VAR __gomp_default_device
+#define GOMP_DYN_VAR __gomp_dyn
+#define GOMP_MAX_ACTIVE_LEVELS_VAR __gomp_max_active_levels
+#define GOMP_BIND_VAR __gomp_bind
+#define GOMP_NTEAMS_VAR __gomp_nteams
 
 /* Miscellaneous functions.  */
 extern void *GOMP_PLUGIN_malloc (size_t) __attribute__ ((malloc));
diff --git a/libgomp/libgomp.h b/libgomp/libgomp.h
index b9e0391..89d1453 100644
--- a/libgomp/libgomp.h
+++ b/libgomp/libgomp.h
@@ -473,6 +473,72 @@ struct gomp_task_icv
   struct target_mem_desc *target_data;
 };
 
+enum gomp_env_var_suffix_t
+{
+  GOMP_ENV_VAR_SUFFIX_UNKNOWN = 0,
+  GOMP_ENV_VAR_SUFFIX_NONE = 1,
+  GOMP_ENV_VAR_SUFFIX_DEV = 2,
+  GOMP_ENV_VAR_SUFFIX_ALL = 4
+};
+
+/* Struct that contains all ICVs for which we need to store initial values.
+   Keeping the initial values is needed for omp_display_env and also used for
+   transmitting device-specific values to the target.  */
+struct gomp_initial_icv_t
+{
+  unsigned long nthreads_var;
+  unsigned long *nthreads_var_list;
+  unsigned long nthreads_var_list_len;
+  enum gomp_schedule_type run_sched_var;
+  int run_sched_chunk_size;
+  int default_device_var;
+  unsigned int thread_limit_var;
+  bool dyn_var;
+  unsigned char max_active_levels_var;
+  char bind_var;
+  char *bind_var_list;
+  unsigned long bind_var_list_len;
+  int nteams_var;
+  int teams_thread_limit_var;
+  int wait_policy;
+  unsigned long stacksize;
+};
+
+struct gomp_icv_flags_t
+{
+  enum gomp_env_var_suffix_t nthreads_var;
+  enum gomp_env_var_suffix_t run_sched_var;
+  enum gomp_env_var_suffix_t run_sched_chunk_size;
+  enum gomp_env_var_suffix_t thread_limit_var;
+  enum gomp_env_var_suffix_t dyn_var;
+  enum gomp_env_var_suffix_t max_active_levels_var;
+  enum gomp_env_var_suffix_t bind_var;
+  enum gomp_env_var_suffix_t nteams_var;
+  enum gomp_env_var_suffix_t stacksize;
+  enum gomp_env_var_suffix_t wait_policy;
+  enum gomp_env_var_suffix_t teams_thread_limit_var;
+};
+
+struct gomp_icv_list {
+  int device_num;
+  void* value;
+  struct gomp_icv_list *next;
+};
+
+extern void *gomp_get_icv_value_ptr (struct gomp_icv_list **list,
+				     int device_num);
+extern struct gomp_icv_list *gomp_run_sched_var_dev_list;
+extern struct gomp_icv_list *gomp_run_sched_chunk_size_dev_list;
+extern struct gomp_icv_list *gomp_nteams_var_dev_list;
+extern struct gomp_icv_list *gomp_max_active_levels_var_dev_list;
+extern struct gomp_icv_list *gomp_proc_bind_var_dev_list;
+extern struct gomp_icv_list *gomp_proc_bind_var_list_dev_list;
+extern struct gomp_icv_list *gomp_proc_bind_var_list_len_dev_list;
+
+extern struct gomp_initial_icv_t gomp_initial_icv_all;
+extern struct gomp_initial_icv_t gomp_initial_icv_dev;
+extern struct gomp_icv_flags_t gomp_initial_icv_flags;
+
 enum gomp_target_offload_t
 {
   GOMP_TARGET_OFFLOAD_DEFAULT,
diff --git a/libgomp/plugin/plugin-gcn.c b/libgomp/plugin/plugin-gcn.c
index f305d72..44ab369 100644
--- a/libgomp/plugin/plugin-gcn.c
+++ b/libgomp/plugin/plugin-gcn.c
@@ -560,6 +560,21 @@ struct heap {
   char data[0];
 };
 
+/* Strings of ICVs which are copied to the device.  */
+static char *GOMP_ICV_STRINGS[] =
+{
+  XSTRING (GOMP_DEVICE_NUM_VAR),
+  XSTRING (GOMP_NTHREADS_VAR),
+  XSTRING (GOMP_THREAD_LIMIT_VAR),
+  XSTRING (GOMP_RUN_SCHED_VAR),
+  XSTRING (GOMP_RUN_SCHED_CHUNK_SIZE),
+  XSTRING (GOMP_DEFAULT_DEVICE_VAR),
+  XSTRING (GOMP_DYN_VAR),
+  XSTRING (GOMP_MAX_ACTIVE_LEVELS_VAR),
+  XSTRING (GOMP_BIND_VAR),
+  XSTRING (GOMP_NTEAMS_VAR)
+};
+
 /* }}}  */
 /* {{{ Global variables  */
 
@@ -3356,7 +3371,7 @@ GOMP_OFFLOAD_load_image (int ord, unsigned version, const void *target_data,
   struct kernel_info *kernel;
   int kernel_count = image_desc->kernel_count;
   unsigned var_count = image_desc->global_variable_count;
-  int other_count = 1;
+  int other_count = 10;
 
   agent = get_agent_info (ord);
   if (!agent)
@@ -3453,36 +3468,43 @@ GOMP_OFFLOAD_load_image (int ord, unsigned version, const void *target_data,
 	}
     }
 
-  GCN_DEBUG ("Looking for variable %s\n", XSTRING (GOMP_DEVICE_NUM_VAR));
-
-  hsa_status_t status;
-  hsa_executable_symbol_t var_symbol;
-  status = hsa_fns.hsa_executable_get_symbol_fn (agent->executable, NULL,
-						 XSTRING (GOMP_DEVICE_NUM_VAR),
-						 agent->id, 0, &var_symbol);
-  if (status == HSA_STATUS_SUCCESS)
+  for (unsigned i = 0; i < other_count; i++)
     {
-      uint64_t device_num_varptr;
-      uint32_t device_num_varsize;
+      GCN_DEBUG ("Looking for variable %s\n", GOMP_ICV_STRINGS[i]);
 
-      status = hsa_fns.hsa_executable_symbol_get_info_fn
-	(var_symbol, HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_ADDRESS,
-	 &device_num_varptr);
-      if (status != HSA_STATUS_SUCCESS)
-	hsa_fatal ("Could not extract a variable from its symbol", status);
-      status = hsa_fns.hsa_executable_symbol_get_info_fn
-	(var_symbol, HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_SIZE,
-	 &device_num_varsize);
-      if (status != HSA_STATUS_SUCCESS)
-	hsa_fatal ("Could not extract a variable size from its symbol", status);
+      hsa_status_t status;
+      hsa_executable_symbol_t var_symbol;
+      status = hsa_fns.hsa_executable_get_symbol_fn (agent->executable, NULL,
+						     GOMP_ICV_STRINGS[i],
+						     agent->id, 0, &var_symbol);
+      if (status == HSA_STATUS_SUCCESS)
+	{
+	  uint64_t varptr;
+	  uint32_t varsize;
 
-      pair->start = device_num_varptr;
-      pair->end = device_num_varptr + device_num_varsize;
+	  status = hsa_fns.hsa_executable_symbol_get_info_fn
+	    (var_symbol, HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_ADDRESS,
+	     &varptr);
+	  if (status != HSA_STATUS_SUCCESS)
+	    hsa_fatal ("Could not extract a variable from its symbol", status);
+	  status = hsa_fns.hsa_executable_symbol_get_info_fn
+	    (var_symbol, HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_SIZE,
+	     &varsize);
+	  if (status != HSA_STATUS_SUCCESS)
+	    hsa_fatal ("Could not extract a variable size from its symbol",
+		       status);
+
+	  pair->start = varptr;
+	  pair->end = varptr + varsize;
+	}
+      else
+	{
+	  /* The variable was not in this image.  */
+	  GCN_DEBUG ("Variable not found in image: %s\n", GOMP_ICV_STRINGS[i]);
+	  pair->start = pair->end = 0;
+	}
+      pair++;
     }
-  else
-    /* The 'GOMP_DEVICE_NUM_VAR' variable was not in this image.  */
-    pair->start = pair->end = 0;
-  pair++;
 
   /* Ensure that constructors are run first.  */
   struct GOMP_kernel_launch_attributes kla =
diff --git a/libgomp/plugin/plugin-nvptx.c b/libgomp/plugin/plugin-nvptx.c
index b4f0a84..5034d5e 100644
--- a/libgomp/plugin/plugin-nvptx.c
+++ b/libgomp/plugin/plugin-nvptx.c
@@ -330,6 +330,21 @@ struct ptx_device
 
 static struct ptx_device **ptx_devices;
 
+/* Strings of ICVs which are copied to the device.  */
+static char *GOMP_ICV_STRINGS[] =
+{
+  XSTRING (GOMP_DEVICE_NUM_VAR),
+  XSTRING (GOMP_NTHREADS_VAR),
+  XSTRING (GOMP_THREAD_LIMIT_VAR),
+  XSTRING (GOMP_RUN_SCHED_VAR),
+  XSTRING (GOMP_RUN_SCHED_CHUNK_SIZE),
+  XSTRING (GOMP_DEFAULT_DEVICE_VAR),
+  XSTRING (GOMP_DYN_VAR),
+  XSTRING (GOMP_MAX_ACTIVE_LEVELS_VAR),
+  XSTRING (GOMP_BIND_VAR),
+  XSTRING (GOMP_NTEAMS_VAR)
+};
+
 static inline struct nvptx_thread *
 nvptx_thread (void)
 {
@@ -1266,7 +1281,7 @@ GOMP_OFFLOAD_load_image (int ord, unsigned version, const void *target_data,
   CUmodule module;
   const char *const *var_names;
   const struct targ_fn_launch *fn_descs;
-  unsigned int fn_entries, var_entries, other_entries, i, j;
+  unsigned int fn_entries, var_entries, other_entries, i, j, k;
   struct targ_fn_descriptor *targ_fns;
   struct addr_pair *targ_tbl;
   const nvptx_tdata_t *img_header = (const nvptx_tdata_t *) target_data;
@@ -1296,8 +1311,8 @@ GOMP_OFFLOAD_load_image (int ord, unsigned version, const void *target_data,
   fn_entries = img_header->fn_num;
   fn_descs = img_header->fn_descs;
 
-  /* Currently, the only other entry kind is 'device number'.  */
-  other_entries = 1;
+  /* Currently, other entry kinds are 'device number' and further ICVs.  */
+  other_entries = 10;
 
   targ_tbl = GOMP_PLUGIN_malloc (sizeof (struct addr_pair)
 				 * (fn_entries + var_entries + other_entries));
@@ -1349,20 +1364,23 @@ GOMP_OFFLOAD_load_image (int ord, unsigned version, const void *target_data,
       targ_tbl->end = targ_tbl->start + bytes;
     }
 
-  CUdeviceptr device_num_varptr;
-  size_t device_num_varsize;
-  CUresult r = CUDA_CALL_NOCHECK (cuModuleGetGlobal, &device_num_varptr,
-				  &device_num_varsize, module,
-				  XSTRING (GOMP_DEVICE_NUM_VAR));
-  if (r == CUDA_SUCCESS)
+  for (k = 0; k < other_entries; k++, targ_tbl++)
     {
-      targ_tbl->start = (uintptr_t) device_num_varptr;
-      targ_tbl->end = (uintptr_t) (device_num_varptr + device_num_varsize);
+      CUdeviceptr varptr;
+      size_t varsize;
+      CUresult r = CUDA_CALL_NOCHECK (cuModuleGetGlobal, &varptr, &varsize,
+				      module, GOMP_ICV_STRINGS[k]);
+      if (r == CUDA_SUCCESS)
+	{
+	  targ_tbl->start = (uintptr_t) varptr;
+	  targ_tbl->end = (uintptr_t) (varptr + varsize);
+	}
+      else
+	{
+	  /* The variable was not in this image.  */
+	  targ_tbl->start = targ_tbl->end = 0;
+	}
     }
-  else
-    /* The 'GOMP_DEVICE_NUM_VAR' variable was not in this image.  */
-    targ_tbl->start = targ_tbl->end = 0;
-  targ_tbl++;
 
   nvptx_set_clocktick (module, dev);
 
diff --git a/libgomp/target.c b/libgomp/target.c
index 698ff14..dd3b7e3 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -2072,6 +2072,21 @@ gomp_update (struct gomp_device_descr *devicep, size_t mapnum, void **hostaddrs,
   gomp_mutex_unlock (&devicep->lock);
 }
 
+/* Helper function for 'gomp_load_image_to_device'. Returns the pointer for an
+   ICV value depending on the device num DEV_NUM and the variable hierarchy
+   (_DEV_42, _DEV, _ALL).  */
+static void*
+gomp_get_icv (struct gomp_icv_list **list, int dev_num,
+	      enum gomp_env_var_suffix_t flag, void *dev_val, void *all_val)
+{
+  void *val = gomp_get_icv_value_ptr (list, dev_num);
+  if (val == NULL && (flag & GOMP_ENV_VAR_SUFFIX_DEV))
+      val = dev_val;
+  if (val == NULL && (flag & GOMP_ENV_VAR_SUFFIX_ALL))
+      val = all_val;
+  return val;
+}
+
 /* Load image pointed by TARGET_DATA to the device, specified by DEVICEP.
    And insert to splay tree the mapping between addresses from HOST_TABLE and
    from loaded target image.  We rely in the host and device compiler
@@ -2093,7 +2108,7 @@ gomp_load_image_to_device (struct gomp_device_descr *devicep, unsigned version,
   int num_vars  = (host_vars_end - host_var_table) / 2;
 
   /* Others currently is only 'device_num' */
-  int num_others = 1;
+  int num_others = 10;
 
   /* Load image to device and get target addresses for the image.  */
   struct addr_pair *target_table = NULL;
@@ -2177,32 +2192,83 @@ gomp_load_image_to_device (struct gomp_device_descr *devicep, unsigned version,
       array++;
     }
 
-  /* Last entry is for the on-device 'device_num' variable. Tolerate case
-     where plugin does not return this entry.  */
+  /* Last entries are for the on-device 'device_num' variable and other ICVs.
+     The position of the variables in TARGET_TABLE results from GOMP_ICV_STRINGS
+     array. Tolerate case where plugin does not return those entries.  */
   if (num_funcs + num_vars < num_target_entries)
     {
-      struct addr_pair *device_num_var = &target_table[num_funcs + num_vars];
-      /* Start address will be non-zero for last entry if GOMP_DEVICE_NUM_VAR
-	 was found in this image.  */
-      if (device_num_var->start != 0)
+      /* The index of the devicep within devices[] is regarded as its
+	 'device number', which is different from the per-device type
+	 devicep->target_id.  */
+      int dev_num = (int) (devicep - &devices[0]);
+      for (i = 0; i < num_others; ++i)
 	{
-	  /* The index of the devicep within devices[] is regarded as its
-	     'device number', which is different from the per-device type
-	     devicep->target_id.  */
-	  int device_num_val = (int) (devicep - &devices[0]);
-	  if (device_num_var->end - device_num_var->start != sizeof (int))
+	  struct addr_pair *var = &target_table[num_funcs + num_vars + i];
+	  /* Start address will be non-zero for the current entry if
+	     the variable was found in this image.  */
+	  if (var->start != 0)
 	    {
-	      gomp_mutex_unlock (&devicep->lock);
-	      if (is_register_lock)
-		gomp_mutex_unlock (&register_lock);
-	      gomp_fatal ("offload plugin managed 'device_num' not of expected "
-			  "format");
-	    }
-
-	  /* Copy device_num value to place on device memory, hereby actually
-	     designating its device number into effect.  */
-	  gomp_copy_host2dev (devicep, NULL, (void *) device_num_var->start,
-			      &device_num_val, sizeof (int), false, NULL);
+	      void *val = NULL;
+	      switch (i)
+		{
+		  case 0:  // GOMP_DEVICE_NUM_VAR
+		    val = &dev_num;
+		    break;
+		  case 1:  // GOMP_NTHREADS_VAR
+		    /* This is ignored since we don't want to overwrite the
+		       values that are set during kernel entering.  */
+		    break;
+		  case 2:  // GOMP_THREAD_LIMIT_VAR
+		    /* This is ignored since we don't want to overwrite the
+		       values that are set during kernel entering.  */
+		    break;
+		  case 3:  // GOMP_RUN_SCHED_VAR
+		    val = gomp_get_icv (&gomp_run_sched_var_dev_list, dev_num,
+					gomp_initial_icv_flags.run_sched_var,
+					&gomp_initial_icv_dev.run_sched_var,
+					&gomp_initial_icv_all.run_sched_var);
+		    break;
+		  case 4:  // GOMP_RUN_SCHED_CHUNK_SIZE
+		    val = gomp_get_icv (&gomp_run_sched_chunk_size_dev_list, dev_num,
+					gomp_initial_icv_flags.run_sched_chunk_size,
+					&gomp_initial_icv_dev.run_sched_chunk_size,
+					&gomp_initial_icv_all.run_sched_chunk_size);
+		    break;
+		  case 5:  // GOMP_DEFAULT_DEVICE_VAR
+		    val = &gomp_global_icv.default_device_var;
+		    break;
+		  case 6:  // GOMP_DYN_VAR
+		    /* This is ignored since we don't want to overwrite the
+		       values that are set during kernel entering.  */
+		    break;
+		  case 7:  // GOMP_MAX_ACTIVE_LEVELS_VAR
+		    val = gomp_get_icv (&gomp_max_active_levels_var_dev_list, dev_num,
+					gomp_initial_icv_flags.max_active_levels_var,
+					&gomp_initial_icv_dev.max_active_levels_var,
+					&gomp_initial_icv_all.max_active_levels_var);
+		    break;
+		  case 8:  // GOMP_BIND_VAR
+		    val = gomp_get_icv (&gomp_proc_bind_var_dev_list, dev_num,
+					gomp_initial_icv_flags.bind_var,
+					&gomp_initial_icv_dev.bind_var,
+					&gomp_initial_icv_all.bind_var);
+		    break;
+		  case 9: // GOMP_NTEAMS_VAR
+		    val = gomp_get_icv (&gomp_nteams_var_dev_list, dev_num,
+					gomp_initial_icv_flags.nteams_var,
+					&gomp_initial_icv_dev.nteams_var,
+					&gomp_initial_icv_all.nteams_var);
+		    break;
+		}
+	      if (val != NULL)
+		{
+		  size_t var_size = var->end - var->start;
+		  /* Copy device_num value to place on device memory, hereby
+		     actually designating its device number into effect.  */
+		  gomp_copy_host2dev (devicep, NULL, (void *) var->start, val,
+				      var_size, false, NULL);
+		}
+	  }
 	}
     }
 
diff --git a/libgomp/testsuite/libgomp.c-c++-common/icv-5.c b/libgomp/testsuite/libgomp.c-c++-common/icv-5.c
new file mode 100644
index 0000000..21bf44d
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/icv-5.c
@@ -0,0 +1,48 @@
+/* { dg-do run } */
+
+#include <omp.h>
+#include <stdlib.h>
+#include <string.h>
+#include <unistd.h>
+
+char const *varnames[] = {
+  "OMP_NUM_TEAMS_DEV_0",
+  "OMP_NUM_TEAMS_DEV_1",
+  "OMP_NUM_TEAMS_DEV_2",
+  "OMP_NUM_TEAMS_ALL",
+  "OMP_NUM_TEAMS_DEV",
+  "OMP_NUM_TEAMS"
+};
+char const *values[] = { "42", "43", "44", "45", "46", "47" };
+const int cnt = 6;
+
+int
+main (int argc, char *const *argv)
+{
+  int updated = 0;
+  for (int i = 0; i < cnt; i++)
+    {
+      if (getenv (varnames[i]) == NULL
+	  || strcmp (getenv (varnames[i]), values[i]) != 0)
+	{
+	  setenv (varnames[i], values[i], 1);
+	  updated = 1;
+	}
+    }
+  if (updated)
+    {
+      execv (argv[0], argv);
+      abort ();
+    }
+
+  if (omp_get_max_teams () != 47)
+    abort ();
+
+  int num_devices = omp_get_num_devices () > 3 ? 3 : omp_get_num_devices ();
+  for (int i=0; i < num_devices; i++)
+    #pragma omp target device (i)
+      if (omp_get_max_teams () != 42 + i)
+	abort ();
+
+  return 0;
+}
\ No newline at end of file
diff --git a/libgomp/testsuite/libgomp.c-c++-common/icv-6.c b/libgomp/testsuite/libgomp.c-c++-common/icv-6.c
new file mode 100644
index 0000000..943147b
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/icv-6.c
@@ -0,0 +1,63 @@
+/* { dg-do run } */
+
+/* This tests the hierarchical usage of ICVs on the device, i.e. if
+   OMP_NUM_TEAMS_DEV_<device_num> is not configured, then the value of
+   OMP_NUM_TEAMS_DEV should be used.  */
+
+#include <omp.h>
+#include <stdlib.h>
+#include <string.h>
+#include <unistd.h>
+
+char const *varnames[] = {
+  "OMP_NUM_TEAMS_ALL",
+  "OMP_NUM_TEAMS_DEV",
+};
+char const *values[] = { "42", "43" };
+const int cnt = 2;
+
+char const *excludes[] = {
+  "OMP_NUM_TEAMS_DEV_0",
+  "OMP_NUM_TEAMS_DEV_1",
+  "OMP_NUM_TEAMS_DEV_2",
+};
+const int cnt_exludes = 3;
+
+int
+main (int argc, char *const *argv)
+{
+  int updated = 0;
+  for (int i = 0; i < cnt; i++)
+    {
+      if (getenv (varnames[i]) == NULL
+	  || strcmp (getenv (varnames[i]), values[i]) != 0)
+	{
+	  setenv (varnames[i], values[i], 1);
+	  updated = 1;
+	}
+    }
+  for (int i = 0; i < cnt_exludes; i++)
+    {
+      if (getenv (excludes[i]) != NULL)
+	{
+	  unsetenv (excludes[i]);
+	  updated = 1;
+	}
+    }
+  if (updated)
+    {
+      execv (argv[0], argv);
+      abort ();
+    }
+
+  if (omp_get_max_teams () != 42)
+    abort ();
+
+  int num_devices = omp_get_num_devices () > 3 ? 3 : omp_get_num_devices ();
+  for (int i=0; i < num_devices; i++)
+    #pragma omp target device (i)
+      if (omp_get_max_teams () != 43)
+	abort ();
+
+  return 0;
+}
\ No newline at end of file
diff --git a/libgomp/testsuite/libgomp.c-c++-common/icv-7.c b/libgomp/testsuite/libgomp.c-c++-common/icv-7.c
new file mode 100644
index 0000000..857d796
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/icv-7.c
@@ -0,0 +1,66 @@
+/* { dg-do run } */
+
+/* This tests the hierarchical usage of ICVs on the host and on devices, i.e. if
+   OMP_NUM_TEAMS_DEV_<device_num>, OMP_NUM_TEAMS_DEV, and
+   OMP_NUM_TEAMS are not configured, then the value of
+   OMP_NUM_TEAMS_ALL should be used for the host as well as for the
+   devices.  */
+
+#include <omp.h>
+#include <stdlib.h>
+#include <string.h>
+#include <unistd.h>
+
+char const *varnames[] = {
+  "OMP_NUM_TEAMS_ALL"
+};
+char const *values[] = { "42" };
+const int cnt = 1;
+
+char const *excludes[] = {
+  "OMP_NUM_TEAMS_DEV_0",
+  "OMP_NUM_TEAMS_DEV_1",
+  "OMP_NUM_TEAMS_DEV_2",
+  "OMP_NUM_TEAMS_DEV",
+  "OMP_NUM_TEAMS"
+};
+const int cnt_exludes = 5;
+
+int
+main (int argc, char *const *argv)
+{
+  int updated = 0;
+  for (int i = 0; i < cnt; i++)
+    {
+      if (getenv (varnames[i]) == NULL
+	  || strcmp (getenv (varnames[i]), values[i]) != 0)
+	{
+	  setenv (varnames[i], values[i], 1);
+	  updated = 1;
+	}
+    }
+  for (int i = 0; i < cnt_exludes; i++)
+    {
+      if (getenv (excludes[i]) != NULL)
+	{
+	  unsetenv (excludes[i]);
+	  updated = 1;
+	}
+    }
+  if (updated)
+    {
+      execv (argv[0], argv);
+      abort ();
+    }
+
+  if (omp_get_max_teams () != 42)
+    abort ();
+
+  int num_devices = omp_get_num_devices () > 3 ? 3 : omp_get_num_devices ();
+  for (int i=0; i < num_devices; i++)
+    #pragma omp target device (i)
+      if (omp_get_max_teams () != 42)
+	abort ();
+
+  return 0;
+}
\ No newline at end of file
diff --git a/libgomp/testsuite/libgomp.c-c++-common/omp-display-env-1.c b/libgomp/testsuite/libgomp.c-c++-common/omp-display-env-1.c
new file mode 100644
index 0000000..0b37e3f
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/omp-display-env-1.c
@@ -0,0 +1,170 @@
+/* { dg-do run } */
+
+#include <omp.h>
+#include <stdlib.h>
+#include <string.h>
+#include <unistd.h>
+
+char const *varnames[] = {
+  "OMP_THREAD_LIMIT_DEV_24",
+  "OMP_THREAD_LIMIT_ALL",
+  "OMP_THREAD_LIMIT_DEV",
+  "OMP_THREAD_LIMIT",
+  "OMP_DEFAULT_DEVICE",
+  "OMP_SCHEDULE_DEV_24",
+  "OMP_SCHEDULE_ALL",
+  "OMP_SCHEDULE_DEV",
+  "OMP_SCHEDULE",
+  "OMP_DYNAMIC_DEV_24",
+  "OMP_DYNAMIC_ALL",
+  "OMP_DYNAMIC_DEV",
+  "OMP_DYNAMIC",
+  "OMP_NUM_THREADS",
+  "OMP_NUM_THREADS_ALL",
+  "OMP_NUM_THREADS_DEV",
+  "OMP_NUM_THREADS_DEV_24",
+  "OMP_MAX_ACTIVE_LEVELS",
+  "OMP_MAX_ACTIVE_LEVELS_ALL",
+  "OMP_MAX_ACTIVE_LEVELS_DEV",
+  "OMP_MAX_ACTIVE_LEVELS_DEV_24",
+  "OMP_NUM_TEAMS",
+  "OMP_NUM_TEAMS_ALL",
+  "OMP_NUM_TEAMS_DEV",
+  "OMP_NUM_TEAMS_DEV_24",
+  "OMP_PROC_BIND",
+  "OMP_PROC_BIND_ALL",
+  "OMP_PROC_BIND_DEV",
+  "OMP_PROC_BIND_DEV_24",
+  "OMP_STACKSIZE",
+  "OMP_STACKSIZE_ALL",
+  "OMP_STACKSIZE_DEV",
+  "OMP_STACKSIZE_DEV_24",
+  "OMP_WAIT_POLICY",
+  "OMP_WAIT_POLICY_ALL",
+  "OMP_WAIT_POLICY_DEV",
+  "OMP_WAIT_POLICY_DEV_24",
+  "OMP_TEAMS_THREAD_LIMIT",
+  "OMP_TEAMS_THREAD_LIMIT_ALL",
+  "OMP_TEAMS_THREAD_LIMIT_DEV",
+  "OMP_TEAMS_THREAD_LIMIT_DEV_24"
+};
+
+char const *values[] = {
+  "42",
+  "43",
+  "44",
+  "45",
+  "42",
+  "guided,4",
+  "dynamic",
+  "guided,1",
+  "guided,2",
+  "true",
+  "true",
+  "true",
+  "true",
+  "4,3,2",
+  "45,46,47",
+  "42,43,44",
+  "14,13,12",
+  "42",
+  "43",
+  "44",
+  "45",
+  "42",
+  "43",
+  "44",
+  "45",
+  "spread",
+  "false",
+  "spread,spread",
+  "spread,close",
+  "42",
+  "42 M",
+  "43 k",
+  "44",
+  "active",
+  "ACTIVE",
+  "passive",
+  "PASSIVE",
+  "42",
+  "43",
+  "44",
+  "45"
+};
+
+const int cnt = 41;
+
+int
+main (int argc, char *const *argv)
+{
+  int updated = 0;
+  for (int i = 0; i < cnt; i++)
+    {
+      if (getenv (varnames[i]) == NULL
+	  || strcmp (getenv (varnames[i]), values[i]) != 0)
+	{
+	  setenv (varnames[i], values[i], 1);
+	  updated = 1;
+	}
+    }
+  if (updated)
+    {
+      execv (argv[0], argv);
+      return 0;
+    }
+  omp_display_env (1);
+  return 0;
+}
+
+/* { dg-output ".*\\\[host] OMP_DYNAMIC = 'TRUE'.*" } */
+/* { dg-output ".*\\\[all\] OMP_DYNAMIC = 'TRUE'.*" } */
+/* { dg-output ".*\\\[device\] OMP_DYNAMIC = 'TRUE'.*" } */
+/* { dg-output ".*\\\[24\] OMP_DYNAMIC = 'TRUE'.*" } */
+
+/* { dg-output ".*\\\[host] OMP_NUM_THREADS = '4,3,2'.*" } */
+/* { dg-output ".*\\\[all\] OMP_NUM_THREADS = '45,46,47'.*" } */
+/* { dg-output ".*\\\[device\] OMP_NUM_THREADS = '42,43,44'.*" } */
+/* { dg-output ".*\\\[24\] OMP_NUM_THREADS = '14,13,12'.*" } */
+
+/* { dg-output ".*\\\[host] OMP_SCHEDULE = 'GUIDED,2'.*" } */
+/* { dg-output ".*\\\[all\] OMP_SCHEDULE = 'DYNAMIC'.*" } */
+/* { dg-output ".*\\\[device\] OMP_SCHEDULE = 'GUIDED'.*" } */
+/* { dg-output ".*\\\[24\] OMP_SCHEDULE = 'GUIDED,4'.*" } */
+
+/* { dg-output ".*\\\[host] OMP_PROC_BIND = 'SPREAD'.*" } */
+/* { dg-output ".*\\\[all\] OMP_PROC_BIND = 'FALSE'.*" } */
+/* { dg-output ".*\\\[device\] OMP_PROC_BIND = 'SPREAD,SPREAD'.*" } */
+/* { dg-output ".*\\\[24\] OMP_PROC_BIND = 'SPREAD,CLOSE'.*" } */
+
+/* { dg-output ".*\\\[host] OMP_STACKSIZE = '43008'.*" } */
+/* { dg-output ".*\\\[all\] OMP_STACKSIZE = '44040192'.*" } */
+/* { dg-output ".*\\\[device\] OMP_STACKSIZE = '44032'.*" } */
+/* { dg-output ".*\\\[24\] OMP_STACKSIZE = '45056'.*" } */
+
+/* { dg-output ".*\\\[host] OMP_WAIT_POLICY = 'ACTIVE'.*" } */
+/* { dg-output ".*\\\[all\] OMP_WAIT_POLICY = 'ACTIVE'.*" } */
+/* { dg-output ".*\\\[device\] OMP_WAIT_POLICY = 'PASSIVE'.*" } */
+/* { dg-output ".*\\\[24\] OMP_WAIT_POLICY = 'PASSIVE'.*" } */
+
+/* { dg-output ".*\\\[host] OMP_THREAD_LIMIT = '45'.*" } */
+/* { dg-output ".*\\\[all\] OMP_THREAD_LIMIT = '43'.*" } */
+/* { dg-output ".*\\\[device\] OMP_THREAD_LIMIT = '44'.*" } */
+/* { dg-output ".*\\\[24\] OMP_THREAD_LIMIT = '42'.*" } */
+
+/* { dg-output ".*\\\[host] OMP_MAX_ACTIVE_LEVELS = '42'.*" } */
+/* { dg-output ".*\\\[all\] OMP_MAX_ACTIVE_LEVELS = '43'.*" } */
+/* { dg-output ".*\\\[device\] OMP_MAX_ACTIVE_LEVELS = '44'.*" } */
+/* { dg-output ".*\\\[24\] OMP_MAX_ACTIVE_LEVELS = '45'.*" } */
+
+/* { dg-output ".*\\\[host] OMP_NUM_TEAMS = '42'.*" } */
+/* { dg-output ".*\\\[all\] OMP_NUM_TEAMS = '43'.*" } */
+/* { dg-output ".*\\\[device\] OMP_NUM_TEAMS = '44'.*" } */
+/* { dg-output ".*\\\[24\] OMP_NUM_TEAMS = '45'.*" } */
+
+/* { dg-output ".*\\\[host] OMP_TEAMS_THREAD_LIMIT = '42'.*" } */
+/* { dg-output ".*\\\[all\] OMP_TEAMS_THREAD_LIMIT = '43'.*" } */
+/* { dg-output ".*\\\[device\] OMP_TEAMS_THREAD_LIMIT = '44'.*" } */
+/* { dg-output ".*\\\[24\] OMP_TEAMS_THREAD_LIMIT = '45'.*" } */
+
+/* { dg-output ".*\\\[all] OMP_DEFAULT_DEVICE = '42'.*" } */
diff --git a/libgomp/testsuite/libgomp.c-c++-common/omp-display-env-2.c b/libgomp/testsuite/libgomp.c-c++-common/omp-display-env-2.c
new file mode 100644
index 0000000..9789430
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/omp-display-env-2.c
@@ -0,0 +1,32 @@
+/* { dg-do run } */
+
+/* This test checks if omp_display_env outputs the initial ICV values.  */
+
+#include <omp.h>
+#include <stdlib.h>
+#include <string.h>
+#include <unistd.h>
+
+int
+main (int argc, char *const *argv)
+{
+  int updated = 0;
+  if (getenv ("OMP_NUM_TEAMS") == NULL
+      || strcmp (getenv ("OMP_NUM_TEAMS"), "42") != 0)
+    {
+      setenv ("OMP_NUM_TEAMS", "42", 1);
+      updated = 1;
+    }
+  if (updated)
+    execv (argv[0], argv);
+
+  omp_display_env (1);
+  omp_set_num_teams (24);
+  if (omp_get_max_teams () != 24)
+    abort ();
+  omp_display_env (1);
+
+  return 0;
+}
+
+/* { dg-output ".*\\\[host] OMP_NUM_TEAMS = '42'.*\\\[host] OMP_NUM_TEAMS = '42'" } */
