@@ -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)
@@ -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)
@@ -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;
@@ -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)
@@ -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)
@@ -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));
@@ -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,
@@ -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 =
@@ -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);
@@ -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 (®ister_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);
+ }
+ }
}
}
new file mode 100644
@@ -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
new file mode 100644
@@ -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
new file mode 100644
@@ -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
new file mode 100644
@@ -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'.*" } */
new file mode 100644
@@ -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'" } */
Hi, This patch considers the environment variable syntax extension for device-specific variants of environment variables from OpenMP 5.1 (see OpenMP 5.1 specification, p. 75 and p. 639). An environment variable (e.g. OMP_NUM_TEAMS) can have different suffixes: _DEV (e.g. OMP_NUM_TEAMS_DEV): affects all devices but not the host. _DEV_<device> (e.g. OMP_NUM_TEAMS_DEV_42): affects only device with number <device>. no suffix (e.g. OMP_NUM_TEAMS): affects only the host. In future OpenMP versions also suffix _ALL will be introduced (see discussion https://github.com/OpenMP/spec/issues/3179). This is also considered in this patch: _ALL (e.g. OMP_NUM_TEAMS_ALL): affects all devices and the host. The precedence is as follows (descending). For the host: 1. no suffix 2. _ALL For devices: 1. _DEV_<device> 2. _DEV 3. _ALL That means, _DEV_<device> is used whenever available. Otherwise _DEV is used if available, and at last _ALL. If there is no value for any of the variable variants, default values are used as already implemented before. This patch concerns parsing (a), storing (b), output (c) and transmission to the device (d): (a) The actual number of devices and the numbering are not known when parsing the environment variables. Thus all environment variables are iterated and searched for device-specific ones. (b) Only configured device-specific variables are stored. Thus, linked lists are used. (c) The output is done in omp_display_env (see specification p. 468f). Global ICVs are tagged with [all], see https://github.com/OpenMP/spec/issues/3179. ICVs which are not global but aren't handled device-specific yet are tagged with [host]. omp_display_env outputs the initial values of the ICVs. That's why separate data structures are introduced (like gomp_initial_icv...). (d) Device-specific ICVs which are already user accessible on the device are transmitted to the device (moreover nteams-var is added and used for the tests). There are ICVs which values are currently set explicitly in the config when copying them to the device: GOMP_NTHREADS_VAR, GOMP_THREAD_LIMIT_VAR, GOMP_DYN_VAR (see gomp_gcn_enter_kernel in libgomp/config/gcn/team.c and gomp_nvptx_main in libgomp/config/nvptx/team.c). The corresponding environment variables are nevertheless parsed and stored device-specific but the transmission to the device is not changed. Marcel ----------------- Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955 OpenMP, libgomp: Environment variable syntax extension. This patch considers the environment variable syntax extension for device-specific variants of environment variables from OpenMP 5.1 (see OpenMP 5.1 specification, p. 75 and p. 639). An environment variable (e.g. OMP_NUM_TEAMS) can have different suffixes: _DEV (e.g. OMP_NUM_TEAMS_DEV): affects all devices but not the host. _DEV_<device> (e.g. OMP_NUM_TEAMS_DEV_42): affects only device with number <device>. no suffix (e.g. OMP_NUM_TEAMS): affects only the host. In future OpenMP versions also suffix _ALL will be introduced (see discussion https://github.com/OpenMP/spec/issues/3179). This is also considered in this patch: _ALL (e.g. OMP_NUM_TEAMS_ALL): affects all devices and the host. The precedence is as follows (descending). For the host: 1. no suffix 2. _ALL For devices: 1. _DEV_<device> 2. _DEV 3. _ALL That means, _DEV_<device> is used whenever available. Otherwise _DEV is used if available, and at last _ALL. If there is no value for any of the variable variants, default values are used as already implemented before. This patch concerns parsing (a), storing (b), output (c) and transmission to the device (d): (a) The actual number of devices and the numbering are not known when parsing the environment variables. Thus all environment variables are iterated and searched for device-specific ones. (b) Only configured device-specific variables are stored. Thus, linked lists are used. (c) The output is done in omp_display_env (see specification p. 468f). Global ICVs are tagged with [all], see https://github.com/OpenMP/spec/issues/3179. ICVs which are not global but aren't handled device-specific yet are tagged with [host]. omp_display_env outputs the initial values of the ICVs. That's why separate data structures are introduced (like gomp_initial_icv...). (d) Device-specific ICVs which are already user accessible on the device are transmitted to the device (moreover nteams-var is added and used for the tests). There are ICVs which values are currently set explicitly in the config when copying them to the device: GOMP_NTHREADS_VAR, GOMP_THREAD_LIMIT_VAR, GOMP_DYN_VAR (see gomp_gcn_enter_kernel in libgomp/config/gcn/team.c and gomp_nvptx_main in libgomp/config/nvptx/team.c). The corresponding environment variables are nevertheless parsed and stored device-specific but the transmission to the device is not changed. libgomp/ChangeLog: * config/gcn/icv-device.c (omp_get_default_device): Return ICV value. (omp_get_max_teams): Added for GCN devices. * config/nvptx/icv-device.c (omp_get_default_device): Return ICV value. (omp_get_max_teams): Added for nvptx devices. * env.c (gomp_get_icv_list): New helper function to return the element of the list of device-specific ICV values for the specified device number. (gomp_get_icv_value_ptr): Like gomp_get_icv_list but returns the actual ICV value. (parse_schedule): Generalized for different variants of OMP_SCHEDULE. (parse_wait_policy): Generalized for different variants of OMP_WAIT_POLICY. (omp_display_env): Extended to output different variants of environment variables. (print_schedule): New helper function for omp_display_env which prints the values of run_sched_var. (print_proc_bind): New helper function for omp_display_env which prints the values of proc_bind_var. (get_device_num): New helper function for parse_device_specific. Extracts the device number from an environment variable name. (add_device_specific_icv): New helper function for parse_device_specific. Adds a new node to the given list. (parse_device_specific): New helper function for 'initialize_env' to parse device-specific environment variables. (initialize_env): Extended to parse the new syntax of environment variables. * icv-device.c (omp_get_max_teams): Added omp_get_max_teams. * icv.c (omp_get_max_teams): Moved to icv-device.c. * libgomp-plugin.h (GOMP_NTHREADS_VAR): New target-side variable that holds the designated ICV of the target device. (GOMP_THREAD_LIMIT_VAR): Likewise. (GOMP_RUN_SCHED_VAR): Likewise. (GOMP_RUN_SCHED_CHUNK_SIZE): Likewise. (GOMP_DEFAULT_DEVICE_VAR): Likewise. (GOMP_DYN_VAR): Likewise. (GOMP_MAX_ACTIVE_LEVELS_VAR): Likewise. (GOMP_BIND_VAR): Likewise. (GOMP_NTEAMS_VAR): Likewise. * libgomp.h (enum gomp_env_var_suffix_t): Enum of flags which are used to store with which suffixes (for which variants) environment variables are defined. (struct gomp_initial_icv_t): Ccontains all ICVs for which we need to store initial values. (struct gomp_icv_flags_t): Store with which suffixes (for which variants) environment variables are defined. (struct gomp_icv_list): Definition of a linked list that is used for storing ICVs for the devices. (gomp_get_icv_value_ptr): Implemented in env.c. * plugin/plugin-gcn.c (GOMP_OFFLOAD_load_image): Extended to read further ICVs from the offload image. * plugin/plugin-nvptx.c (GOMP_OFFLOAD_load_image): Extended to read further ICVs from the offload image. * target.c (gomp_get_icv): New helper function for 'gomp_load_image_to_device' that returns the pointer for an ICV value depending on the device number and the variable hierarchy. (gomp_load_image_to_device): Extended to copy further ICVs to the device. * testsuite/libgomp.c-c++-common/icv-5.c: New test. * testsuite/libgomp.c-c++-common/icv-6.c: New test. * testsuite/libgomp.c-c++-common/icv-7.c: New test. * testsuite/libgomp.c-c++-common/omp-display-env-1.c: New test. * testsuite/libgomp.c-c++-common/omp-display-env-2.c: New test.