mirror/gcc
2
0
Fork 0
mirror of git://gcc.gnu.org/git/gcc.git synced 2025-04-13 14:51:18 +08:00
Code Issues Packages Projects Releases Wiki Activity

OpenMP, libgomp: Environment variable syntax extension

Browse Source 
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, a linked list
is 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 is
why a dedicated data structure is introduced for the inital values only
(gomp_initial_icv_list).
(d) Device-specific ICVs are transmitted to the device via GOMP_ADDITIONAL_ICVS.

libgomp/ChangeLog:

	* config/gcn/icv-device.c (omp_get_default_device): Return device-
	specific ICV.
	(omp_get_max_teams): Added for GCN devices.
	(omp_set_num_teams): Likewise.
	(ialias): Likewise.
	* config/nvptx/icv-device.c (omp_get_default_device): Return device-
	specific ICV.
	(omp_get_max_teams): Added for NVPTX devices.
	(omp_set_num_teams): Likewise.
	(ialias): Likewise.
	* env.c (struct gomp_icv_list): New struct to store entries of initial
	ICV values.
	(struct gomp_offload_icv_list): New struct to store entries of device-
	specific ICV values that are copied to the device and back.
	(struct gomp_default_icv_values): New struct to store default values of
	ICVs according to the OpenMP standard.
	(parse_schedule): Generalized for different variants of OMP_SCHEDULE.
	(print_env_var_error): Function that prints an error for invalid values
	for ICVs.
	(parse_unsigned_long_1): Removed getenv.  Generalized.
	(parse_unsigned_long): Likewise.
	(parse_int_1): Likewise.
	(parse_int): Likewise.
	(parse_int_secure): Likewise.
	(parse_unsigned_long_list): Likewise.
	(parse_target_offload): Likewise.
	(parse_bind_var): Likewise.
	(parse_stacksize): Likewise.
	(parse_boolean): Likewise.
	(parse_wait_policy): Likewise.
	(parse_allocator): Likewise.
	(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.
	(enum gomp_parse_type): Collection of types used for parsing environment
	variables.
	(ENTRY): Preprocess string lengths of environment variables.
	(OMP_VAR_CNT): Preprocess table size.
	(OMP_HOST_VAR_CNT): Likewise.
	(INT_MAX_STR_LEN): Constant for the maximal number of digits of a device
	number.
	(gomp_get_icv_flag): Returns if a flag for a particular ICV is set.
	(gomp_set_icv_flag): Sets a flag for a particular ICV.
	(print_device_specific_icvs): New helper function for omp_display_env to
	print device specific ICV values.
	(get_device_num): New helper function for parse_device_specific.
	Extracts the device number from an environment variable name.
	(get_icv_member_addr): Gets the memory address for a particular member
	of an ICV struct.
	(gomp_get_initial_icv_item): Get a list item of gomp_initial_icv_list.
	(initialize_icvs): New function to initialize a gomp_initial_icvs
	struct.
	(add_initial_icv_to_list): Adds an ICV struct to gomp_initial_icv_list.
	(startswith): Checks if a string starts with a given prefix.
	(initialize_env): Extended to parse the new syntax of environment
	variables.
	* icv-device.c (omp_get_max_teams): Added.
	(ialias): Likewise.
	(omp_set_num_teams): Likewise.
	* icv.c (omp_set_num_teams): Moved to icv-device.c.
	(omp_get_max_teams): Likewise.
	(ialias): Likewise.
	* libgomp-plugin.h (GOMP_DEVICE_NUM_VAR): Removed.
	(GOMP_ADDITIONAL_ICVS): New target-side struct that
	holds the designated ICVs of the target device.
	* libgomp.h (enum gomp_icvs): Collection of ICVs.
	(enum gomp_device_num): Definition of device numbers for _ALL, _DEV, and
	no suffix.
	(enum gomp_env_suffix): Collection of possible suffixes of environment
	variables.
	(struct gomp_initial_icvs): Contains all ICVs for which we need to store
	initial values.
	(struct gomp_default_icv):New struct to hold ICVs for which we need
	to store initial values.
	(struct gomp_icv_list): Definition of a linked list that is used for
	storing ICVs for the devices and also for _DEV, _ALL, and without
	suffix.
	(struct gomp_offload_icvs): New struct to hold ICVs that are copied to
	a device.
	(struct gomp_offload_icv_list): Definition of a linked list that holds
	device-specific ICVs that are copied to devices.
	(gomp_get_initial_icv_item): Get a list item of gomp_initial_icv_list.
	(gomp_get_icv_flag): Returns if a flag for a particular ICV is set.
	* libgomp.texi: Updated.
	* 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): Likewise.
	* target.c (gomp_get_offload_icv_item): Get a list item of
	gomp_offload_icv_list.
	(get_gomp_offload_icvs): New. Returns the ICV values
	depending on the device num and the variable hierarchy.
	(gomp_load_image_to_device): Extended to copy further ICVs to a 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/icv-8.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.
This commit is contained in:
Marcel Vollweiler 2022-09-08 10:01:33 -07:00
parent 338a5b0d7d
commit 9f2fca5659
17 changed files with 1580 additions and 380 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

27
libgomp/config/gcn/icv-device.c
View File

@ -28,6 +28,10 @@
#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 struct gomp_offload_icvs GOMP_ADDITIONAL_ICVS;
void
omp_set_default_device (int device_num __attribute__((unused)))
{
@ -36,7 +40,7 @@ omp_set_default_device (int device_num __attribute__((unused)))
int
omp_get_default_device (void)
{
return 0;
return GOMP_ADDITIONAL_ICVS.default_device;
}
int
@ -58,14 +62,23 @@ 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. */
volatile int GOMP_DEVICE_NUM_VAR;
int
omp_get_device_num (void)
{
return GOMP_DEVICE_NUM_VAR;
return GOMP_ADDITIONAL_ICVS.device_num;
}
int
omp_get_max_teams (void)
{
return GOMP_ADDITIONAL_ICVS.nteams;
}
void
omp_set_num_teams (int num_teams)
{
if (num_teams >= 0)
GOMP_ADDITIONAL_ICVS.nteams = num_teams;
}
ialias (omp_set_default_device)
@ -74,3 +87,5 @@ 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)
ialias (omp_set_num_teams)

27
libgomp/config/nvptx/icv-device.c
View File

@ -28,6 +28,10 @@
#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 struct gomp_offload_icvs GOMP_ADDITIONAL_ICVS;
void
omp_set_default_device (int device_num __attribute__((unused)))
{
@ -36,7 +40,7 @@ omp_set_default_device (int device_num __attribute__((unused)))
int
omp_get_default_device (void)
{
return 0;
return GOMP_ADDITIONAL_ICVS.default_device;
}
int
@ -58,14 +62,23 @@ 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;
return GOMP_ADDITIONAL_ICVS.device_num;
}
int
omp_get_max_teams (void)
{
return GOMP_ADDITIONAL_ICVS.nteams;
}
void
omp_set_num_teams (int num_teams)
{
if (num_teams >= 0)
GOMP_ADDITIONAL_ICVS.nteams = num_teams;
}
ialias (omp_set_default_device)
@ -74,3 +87,5 @@ 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)
ialias (omp_set_num_teams)

1318
libgomp/env.c
View File

File diff suppressed because it is too large Load Diff

17
libgomp/icv-device.c
View File

@ -80,3 +80,20 @@ 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)
void
omp_set_num_teams (int num_teams)
{
if (num_teams >= 0)
gomp_nteams_var = num_teams;
}
ialias (omp_set_num_teams)

15
libgomp/icv.c
View File

@ -148,19 +148,6 @@ omp_get_supported_active_levels (void)
return gomp_supported_active_levels;
}
void
omp_set_num_teams (int num_teams)
{
if (num_teams >= 0)
gomp_nteams_var = num_teams;
}
int
omp_get_max_teams (void)
{
return gomp_nteams_var;
}
void
omp_set_teams_thread_limit (int thread_limit)
{
@ -274,8 +261,6 @@ ialias (omp_get_thread_limit)
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)

10
libgomp/libgomp-plugin.h
View File

@ -102,11 +102,11 @@ 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
libgomp code and the offload plugin (which in the latter case must be
stringified). */
#define GOMP_DEVICE_NUM_VAR __gomp_device_num
/* This following symbol is used to name the target side variable struct that
holds the designated ICVs of the target device. The symbol needs to be
available to libgomp code and the offload plugin (which in the latter case
must be stringified). */
#define GOMP_ADDITIONAL_ICVS __gomp_additional_icvs
/* Miscellaneous functions. */
extern void *GOMP_PLUGIN_malloc (size_t) __attribute__ ((malloc));

114
libgomp/libgomp.h
View File

@ -453,6 +453,38 @@ struct gomp_team_state
struct target_mem_desc;
enum gomp_icvs
{
GOMP_ICV_NTEAMS = 1,
GOMP_ICV_SCHEDULE = 2,
GOMP_ICV_SCHEDULE_CHUNK_SIZE = 3,
GOMP_ICV_DYNAMIC = 4,
GOMP_ICV_TEAMS_THREAD_LIMIT = 5,
GOMP_ICV_THREAD_LIMIT = 6,
GOMP_ICV_NTHREADS = 7,
GOMP_ICV_NTHREADS_LIST = 8,
GOMP_ICV_NTHREADS_LIST_LEN = 9,
GOMP_ICV_BIND = 10,
GOMP_ICV_BIND_LIST = 11,
GOMP_ICV_BIND_LIST_LEN = 12,
GOMP_ICV_MAX_ACTIVE_LEVELS = 13,
GOMP_ICV_WAIT_POLICY = 14,
GOMP_ICV_STACKSIZE = 15,
GOMP_ICV_DEFAULT_DEVICE = 16,
GOMP_ICV_CANCELLATION = 17,
GOMP_ICV_DISPLAY_AFFINITY = 18,
GOMP_ICV_TARGET_OFFLOAD = 19,
GOMP_ICV_MAX_TASK_PRIORITY = 20,
GOMP_ICV_ALLOCATOR = 21
};
enum gomp_device_num
{
GOMP_DEVICE_NUM_FOR_DEV = -1,
GOMP_DEVICE_NUM_FOR_ALL = -2,
GOMP_DEVICE_NUM_FOR_NO_SUFFIX = -3
};
/* These are the OpenMP 4.0 Internal Control Variables described in
section 2.3.1. Those described as having one copy per task are
stored within the structure; those described as having one copy
@ -472,6 +504,80 @@ struct gomp_task_icv
struct target_mem_desc *target_data;
};
enum gomp_env_suffix
{
GOMP_ENV_SUFFIX_UNKNOWN = 0,
GOMP_ENV_SUFFIX_NONE = 1,
GOMP_ENV_SUFFIX_DEV = 2,
GOMP_ENV_SUFFIX_ALL = 4,
GOMP_ENV_SUFFIX_DEV_X = 8
};
/* Struct that contains all ICVs for which we need to store initial values.
Keeping the initial values is needed for omp_display_env. Moreover initial
_DEV and _ALL variants of environment variables are also used to determine
actually used values for devices and for the host. */
struct gomp_initial_icvs
{
unsigned long *nthreads_var_list;
char *bind_var_list;
unsigned long nthreads_var;
unsigned long nthreads_var_list_len;
unsigned long bind_var_list_len;
unsigned long stacksize;
int run_sched_chunk_size;
int default_device_var;
int nteams_var;
int teams_thread_limit_var;
int wait_policy;
unsigned int thread_limit_var;
enum gomp_schedule_type run_sched_var;
bool dyn_var;
unsigned char max_active_levels_var;
char bind_var;
};
struct gomp_default_icv
{
unsigned long nthreads_var;
enum gomp_schedule_type run_sched_var;
int run_sched_chunk_size;
int default_device_var;
unsigned int thread_limit_var;
int nteams_var;
int teams_thread_limit_var;
bool dyn_var;
unsigned char max_active_levels_var;
char bind_var;
};
/* DEVICE_NUM "-1" is reserved for "_DEV" icvs.
DEVICE_NUM "-2" is reserved for "_ALL" icvs.
DEVICE_NUM "-3" is reserved for ICVs without suffix.
Non-negative DEVICE_NUM is for "_DEV_X" icvs. */
struct gomp_icv_list
{
int device_num;
uint32_t flags;
struct gomp_initial_icvs icvs;
struct gomp_icv_list *next;
};
struct gomp_offload_icvs
{
int device_num;
int default_device;
int nteams;
int teams_thread_limit;
};
struct gomp_offload_icv_list
{
int device_num;
struct gomp_offload_icvs icvs;
struct gomp_offload_icv_list *next;
};
enum gomp_target_offload_t
{
GOMP_TARGET_OFFLOAD_DEFAULT,
@ -503,6 +609,9 @@ extern bool gomp_display_affinity_var;
extern char *gomp_affinity_format_var;
extern size_t gomp_affinity_format_len;
extern uintptr_t gomp_def_allocator;
extern const struct gomp_default_icv gomp_default_icv_values;
extern struct gomp_icv_list *gomp_initial_icv_list;
extern struct gomp_offload_icv_list *gomp_offload_icv_list;
extern int goacc_device_num;
extern char *goacc_device_type;
extern int goacc_default_dims[GOMP_DIM_MAX];
@ -927,6 +1036,11 @@ extern void gomp_display_affinity_thread (gomp_thread_handle,
struct gomp_team_state *,
unsigned int) __attribute__((cold));
/* env.c */
extern struct gomp_icv_list *gomp_get_initial_icv_item (int dev_num);
extern bool gomp_get_icv_flag (uint32_t value, enum gomp_icvs icv);
/* iter.c */
extern int gomp_iter_static_next (long *, long *);

2
libgomp/libgomp.texi
View File

@ -284,7 +284,7 @@ The OpenMP 4.5 specification is fully supported.
@item @code{declare variant}: new clauses @code{adjust_args} and
@code{append_args} @tab N @tab
@item @code{dispatch} construct @tab N @tab
@item device-specific ICV settings the environment variables @tab N @tab
@item device-specific ICV settings with environment variables @tab Y @tab
@item @code{assume} directive @tab N @tab
@item @code{nothing} directive @tab Y @tab
@item @code{error} directive @tab Y @tab

29
libgomp/plugin/plugin-gcn.c
View File

@ -3367,6 +3367,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;
/* Currently, "others" is a struct of ICVS. */
int other_count = 1;
agent = get_agent_info (ord);
@ -3464,36 +3465,40 @@ GOMP_OFFLOAD_load_image (int ord, unsigned version, const void *target_data,
}
}
GCN_DEBUG ("Looking for variable %s\n", XSTRING (GOMP_DEVICE_NUM_VAR));
GCN_DEBUG ("Looking for variable %s\n", XSTRING (GOMP_ADDITIONAL_ICVS));
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),
XSTRING (GOMP_ADDITIONAL_ICVS),
agent->id, 0, &var_symbol);
if (status == HSA_STATUS_SUCCESS)
{
uint64_t device_num_varptr;
uint32_t device_num_varsize;
uint64_t varptr;
uint32_t varsize;
status = hsa_fns.hsa_executable_symbol_get_info_fn
(var_symbol, HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_ADDRESS,
&device_num_varptr);
&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);
&varsize);
if (status != HSA_STATUS_SUCCESS)
hsa_fatal ("Could not extract a variable size from its symbol", status);
hsa_fatal ("Could not extract a variable size from its symbol",
status);
pair->start = device_num_varptr;
pair->end = device_num_varptr + device_num_varsize;
pair->start = varptr;
pair->end = varptr + varsize;
}
else
/* The 'GOMP_DEVICE_NUM_VAR' variable was not in this image. */
pair->start = pair->end = 0;
pair++;
{
/* The variable was not in this image. */
GCN_DEBUG ("Variable not found in image: %s\n",
XSTRING (GOMP_ADDITIONAL_ICVS));
pair->start = pair->end = 0;
}
/* Ensure that constructors are run first. */
struct GOMP_kernel_launch_attributes kla =

19
libgomp/plugin/plugin-nvptx.c
View File

@ -1305,7 +1305,7 @@ 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'. */
/* Currently, other_entries contains only the struct of ICVs. */
other_entries = 1;
targ_tbl = GOMP_PLUGIN_malloc (sizeof (struct addr_pair)
@ -1358,20 +1358,19 @@ 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));
CUdeviceptr varptr;
size_t varsize;
CUresult r = CUDA_CALL_NOCHECK (cuModuleGetGlobal, &varptr, &varsize,
module, XSTRING (GOMP_ADDITIONAL_ICVS));
if (r == CUDA_SUCCESS)
{
targ_tbl->start = (uintptr_t) device_num_varptr;
targ_tbl->end = (uintptr_t) (device_num_varptr + device_num_varsize);
targ_tbl->start = (uintptr_t) varptr;
targ_tbl->end = (uintptr_t) (varptr + varsize);
}
else
/* The 'GOMP_DEVICE_NUM_VAR' variable was not in this image. */
/* The variable was not in this image. */
targ_tbl->start = targ_tbl->end = 0;
targ_tbl++;
nvptx_set_clocktick (module, dev);

119
libgomp/target.c
View File

@ -2108,6 +2108,68 @@ gomp_update (struct gomp_device_descr *devicep, size_t mapnum, void **hostaddrs,
gomp_mutex_unlock (&devicep->lock);
}
static struct gomp_offload_icv_list *
gomp_get_offload_icv_item (int dev_num)
{
struct gomp_offload_icv_list *l = gomp_offload_icv_list;
while (l != NULL && l->device_num != dev_num)
l = l->next;
return l;
}
/* Helper function for 'gomp_load_image_to_device'. Returns the ICV values
depending on the device num and the variable hierarchy
(_DEV_42, _DEV, _ALL). If no ICV was initially configured for the given
device and thus no item with that device number is contained in
gomp_offload_icv_list, then a new item is created and added to the list. */
static struct gomp_offload_icvs *
get_gomp_offload_icvs (int dev_num)
{
struct gomp_icv_list *dev
= gomp_get_initial_icv_item (GOMP_DEVICE_NUM_FOR_DEV);
struct gomp_icv_list *all
= gomp_get_initial_icv_item (GOMP_DEVICE_NUM_FOR_ALL);
struct gomp_icv_list *dev_x = gomp_get_initial_icv_item (dev_num);
struct gomp_offload_icv_list *offload_icvs
= gomp_get_offload_icv_item (dev_num);
if (offload_icvs != NULL)
return &offload_icvs->icvs;
struct gomp_offload_icv_list *new
= (struct gomp_offload_icv_list *) gomp_malloc (sizeof (struct gomp_offload_icv_list));
new->device_num = dev_num;
new->icvs.device_num = dev_num;
new->next = gomp_offload_icv_list;
if (dev_x != NULL && gomp_get_icv_flag (dev_x->flags, GOMP_ICV_NTEAMS))
new->icvs.nteams = dev_x->icvs.nteams_var;
else if (dev != NULL && gomp_get_icv_flag (dev->flags, GOMP_ICV_NTEAMS))
new->icvs.nteams = dev->icvs.nteams_var;
else if (all != NULL && gomp_get_icv_flag (all->flags, GOMP_ICV_NTEAMS))
new->icvs.nteams = all->icvs.nteams_var;
else
new->icvs.nteams = gomp_default_icv_values.nteams_var;
if (dev_x != NULL
&& gomp_get_icv_flag (dev_x->flags, GOMP_ICV_DEFAULT_DEVICE))
new->icvs.default_device = dev_x->icvs.default_device_var;
else if (dev != NULL
&& gomp_get_icv_flag (dev->flags, GOMP_ICV_DEFAULT_DEVICE))
new->icvs.default_device = dev->icvs.default_device_var;
else if (all != NULL
&& gomp_get_icv_flag (all->flags, GOMP_ICV_DEFAULT_DEVICE))
new->icvs.default_device = all->icvs.default_device_var;
else
new->icvs.default_device = gomp_default_icv_values.default_device_var;
gomp_offload_icv_list = new;
return &new->icvs;
}
/* 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
@ -2128,9 +2190,6 @@ gomp_load_image_to_device (struct gomp_device_descr *devicep, unsigned version,
int num_funcs = host_funcs_end - host_func_table;
int num_vars = (host_vars_end - host_var_table) / 2;
/* Others currently is only 'device_num' */
int num_others = 1;
/* Load image to device and get target addresses for the image. */
struct addr_pair *target_table = NULL;
int i, num_target_entries;
@ -2140,8 +2199,8 @@ gomp_load_image_to_device (struct gomp_device_descr *devicep, unsigned version,
target_data, &target_table);
if (num_target_entries != num_funcs + num_vars
/* Others (device_num) are included as trailing entries in pair list. */
&& num_target_entries != num_funcs + num_vars + num_others)
/* "+1" due to the additional ICV struct. */
&& num_target_entries != num_funcs + num_vars + 1)
{
gomp_mutex_unlock (&devicep->lock);
if (is_register_lock)
@ -2153,7 +2212,9 @@ gomp_load_image_to_device (struct gomp_device_descr *devicep, unsigned version,
/* Insert host-target address mapping into splay tree. */
struct target_mem_desc *tgt = gomp_malloc (sizeof (*tgt));
tgt->array = gomp_malloc ((num_funcs + num_vars) * sizeof (*tgt->array));
/* "+1" due to the additional ICV struct. */
tgt->array = gomp_malloc ((num_funcs + num_vars + 1)
* sizeof (*tgt->array));
tgt->refcount = REFCOUNT_INFINITY;
tgt->tgt_start = 0;
tgt->tgt_end = 0;
@ -2213,32 +2274,40 @@ 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 entry is for a ICVs variable.
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)
struct addr_pair *var = &target_table[num_funcs + num_vars];
/* Start address will be non-zero for the ICVs variable if
the variable was found in this image. */
if (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 device_num_val = (int) (devicep - &devices[0]);
if (device_num_var->end - device_num_var->start != sizeof (int))
{
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");
}
int dev_num = (int) (devicep - &devices[0]);
struct gomp_offload_icvs *icvs = get_gomp_offload_icvs (dev_num);
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 *) device_num_var->start,
&device_num_val, sizeof (int), false, NULL);
/* Copy the ICVs variable to place on device memory, hereby
actually designating its device number into effect. */
gomp_copy_host2dev (devicep, NULL, (void *) var->start, icvs,
var_size, false, NULL);
splay_tree_key k = &array->key;
k->host_start = (uintptr_t) icvs;
k->host_end =
k->host_start + (size_mask & sizeof (struct gomp_offload_icvs));
k->tgt = tgt;
k->tgt_offset = var->start;
k->refcount = REFCOUNT_INFINITY;
k->dynamic_refcount = 0;
k->aux = NULL;
array->left = NULL;
array->right = NULL;
splay_tree_insert (&devicep->mem_map, array);
array++;
}
}

25
libgomp/testsuite/libgomp.c-c++-common/icv-5.c Normal file
View File

@ -0,0 +1,25 @@
/* { dg-do run } */
/* { dg-set-target-env-var OMP_NUM_TEAMS_DEV_0 "42" } */
/* { dg-set-target-env-var OMP_NUM_TEAMS_DEV_1 "43" } */
/* { dg-set-target-env-var OMP_NUM_TEAMS_DEV_2 "44" } */
/* { dg-set-target-env-var OMP_NUM_TEAMS_ALL "45" } */
/* { dg-set-target-env-var OMP_NUM_TEAMS_DEV "46" } */
/* { dg-set-target-env-var OMP_NUM_TEAMS "47" } */
#include <omp.h>
#include <stdlib.h>
int
main ()
{
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;
}

45
libgomp/testsuite/libgomp.c-c++-common/icv-6.c Normal file
View File

@ -0,0 +1,45 @@
/* { dg-do run } */
/* { dg-set-target-env-var OMP_NUM_TEAMS_ALL "42" } */
/* { dg-set-target-env-var OMP_NUM_TEAMS_DEV "43" } */
/* { dg-set-target-env-var OMP_SCHEDULE_ALL "guided,4" } */
/* { dg-set-target-env-var OMP_DYNAMIC_ALL "true" } */
/* { dg-set-target-env-var OMP_TEAMS_THREAD_LIMIT_ALL "44" } */
/* { dg-set-target-env-var OMP_THREAD_LIMIT_ALL "45" } */
/* { dg-set-target-env-var OMP_NUM_THREADS_ALL "46,3,2" } */
/* { dg-set-target-env-var OMP_MAX_ACTIVE_LEVELS_ALL "47" } */
/* { dg-set-target-env-var OMP_PROC_BIND_ALL "spread" } */
/* { dg-set-target-env-var OMP_WAIT_POLICY_ALL "active" } */
/* 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. And if there is no environment variable
without suffix, then the corresponding _ALL variant should be used. */
#include <omp.h>
#include <stdlib.h>
int
main ()
{
enum omp_sched_t kind;
int chunk_size;
omp_get_schedule(&kind, &chunk_size);
if (omp_get_max_teams () != 42
|| !omp_get_dynamic ()
|| kind != 3 || chunk_size != 4
|| omp_get_teams_thread_limit () != 44
|| omp_get_thread_limit () != 45
|| omp_get_max_threads () != 46
|| omp_get_proc_bind () != omp_proc_bind_spread
|| omp_get_max_active_levels () != 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 () != 43)
abort ();
return 0;
}

26
libgomp/testsuite/libgomp.c-c++-common/icv-7.c Normal file
View File

@ -0,0 +1,26 @@
/* { dg-do run } */
/* { dg-set-target-env-var OMP_NUM_TEAMS_ALL "42" } */
/* 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>
int
main ()
{
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;
}

26
libgomp/testsuite/libgomp.c-c++-common/icv-8.c Normal file
View File

@ -0,0 +1,26 @@
/* { dg-do run } */
/* { dg-set-target-env-var OMP_NUM_TEAMS_DEV_1234567890 "42" } */
/* { dg-set-target-env-var OMP_NUM_TEAMS_DEV_ "43" } */
/* { dg-set-target-env-var OMP_NUM_TEAMS_DEV_01 "44" } */
/* { dg-set-target-env-var OMP_NUM_TEAMS_DEV_a "45" } */
/* { dg-set-target-env-var OMP_NUM_TEAMS_DEV_12345678901 "46" } */
/* { dg-set-target-env-var OMP_NUM_TEAMS_DEV_-1 "47" } */
/* { dg-set-target-env-var "OMP_NUM_TEAMS_DEV_ 1" "48" } */
/* { dg-set-target-env-var "OMP_NUM_TEAMS_DEV_00" "49" } */
#include <omp.h>
#include <stdlib.h>
int
main ()
{
return 0;
}
/* { dg-output ".*Invalid device number in OMP_NUM_TEAMS_DEV_=43.*" { target native } } */
/* { dg-output ".*Invalid device number in OMP_NUM_TEAMS_DEV_01=44.*" { target native } } */
/* { dg-output ".*Invalid device number in OMP_NUM_TEAMS_DEV_a=45.*" { target native } } */
/* { dg-output ".*Invalid device number in OMP_NUM_TEAMS_DEV_12345678901=46.*" { target native } } */
/* { dg-output ".*Invalid device number in OMP_NUM_TEAMS_DEV_-1=47.*" { target native } } */
/* { dg-output ".*Invalid device number in OMP_NUM_TEAMS_DEV_ 1=48.*" { target native } } */
/* { dg-output ".*Invalid device number in OMP_NUM_TEAMS_DEV_00=49.*" { target native } } */

119
libgomp/testsuite/libgomp.c-c++-common/omp-display-env-1.c Normal file
View File

@ -0,0 +1,119 @@
/* { dg-do run } */
/* { dg-set-target-env-var OMP_THREAD_LIMIT_DEV_24 "42" } */
/* { dg-set-target-env-var OMP_THREAD_LIMIT_ALL "43" } */
/* { dg-set-target-env-var OMP_THREAD_LIMIT_DEV "44" } */
/* { dg-set-target-env-var OMP_THREAD_LIMIT "45" } */
/* { dg-set-target-env-var OMP_DEFAULT_DEVICE "42" } */
/* { dg-set-target-env-var OMP_SCHEDULE_DEV_24 "guided,4" } */
/* { dg-set-target-env-var OMP_SCHEDULE_ALL "dynamic" } */
/* { dg-set-target-env-var OMP_SCHEDULE_DEV "guided,1" } */
/* { dg-set-target-env-var OMP_SCHEDULE "guided,2" } */
/* { dg-set-target-env-var OMP_DYNAMIC_DEV_24 "true" } */
/* { dg-set-target-env-var OMP_DYNAMIC_ALL "true" } */
/* { dg-set-target-env-var OMP_DYNAMIC_DEV "true" } */
/* { dg-set-target-env-var OMP_DYNAMIC "true" } */
/* { dg-set-target-env-var OMP_NUM_THREADS "4,3,2" } */
/* { dg-set-target-env-var OMP_NUM_THREADS_ALL "45,46,47" } */
/* { dg-set-target-env-var OMP_NUM_THREADS_DEV "42,43,44" } */
/* { dg-set-target-env-var OMP_NUM_THREADS_DEV_24 "14,13,12" } */
/* { dg-set-target-env-var OMP_MAX_ACTIVE_LEVELS "42" } */
/* { dg-set-target-env-var OMP_MAX_ACTIVE_LEVELS_ALL "43" } */
/* { dg-set-target-env-var OMP_MAX_ACTIVE_LEVELS_DEV "44" } */
/* { dg-set-target-env-var OMP_MAX_ACTIVE_LEVELS_DEV_24 "45" } */
/* { dg-set-target-env-var OMP_NUM_TEAMS "42" } */
/* { dg-set-target-env-var OMP_NUM_TEAMS_ALL "43" } */
/* { dg-set-target-env-var OMP_NUM_TEAMS_DEV "44" } */
/* { dg-set-target-env-var OMP_NUM_TEAMS_DEV_24 "45" } */
/* { dg-set-target-env-var OMP_PROC_BIND "spread" } */
/* { dg-set-target-env-var OMP_PROC_BIND_ALL "close" } */
/* { dg-set-target-env-var OMP_PROC_BIND_DEV "spread,spread" } */
/* { dg-set-target-env-var OMP_PROC_BIND_DEV_24 "spread,close" } */
/* { dg-set-target-env-var OMP_STACKSIZE "42" } */
/* { dg-set-target-env-var OMP_STACKSIZE_ALL "42 M" } */
/* { dg-set-target-env-var OMP_STACKSIZE_DEV "43 k" } */
/* { dg-set-target-env-var OMP_STACKSIZE_DEV_24 "44" } */
/* { dg-set-target-env-var OMP_WAIT_POLICY "active" } */
/* { dg-set-target-env-var OMP_WAIT_POLICY_ALL "ACTIVE" } */
/* { dg-set-target-env-var OMP_WAIT_POLICY_DEV "passive" } */
/* { dg-set-target-env-var OMP_WAIT_POLICY_DEV_24 "PASSIVE" } */
/* { dg-set-target-env-var OMP_TEAMS_THREAD_LIMIT "42" } */
/* { dg-set-target-env-var OMP_TEAMS_THREAD_LIMIT_ALL "43" } */
/* { dg-set-target-env-var OMP_TEAMS_THREAD_LIMIT_DEV "44" } */
/* { dg-set-target-env-var OMP_TEAMS_THREAD_LIMIT_DEV_24 "45" } */
/* { dg-set-target-env-var OMP_CANCELLATION "true" } */
/* { dg-set-target-env-var OMP_DISPLAY_AFFINITY "true" } */
/* { dg-set-target-env-var OMP_TARGET_OFFLOAD "mandatory" } */
/* { dg-set-target-env-var OMP_MAX_TASK_PRIORITY "20" } */
/* { dg-set-target-env-var OMP_ALLOCATOR "omp_const_mem_alloc" } */
/* { dg-set-target-env-var OMP_NESTED "false" } */
#include <omp.h>
#include <stdlib.h>
int
main ()
{
omp_display_env (1);
return 0;
}
/* { dg-output ".*\\\[host] OMP_DYNAMIC = 'TRUE'.*" { target native } } */
/* { dg-output ".*\\\[all] OMP_DYNAMIC = 'TRUE'.*" { target native } } */
/* { dg-output ".*\\\[device] OMP_DYNAMIC = 'TRUE'.*" { target native } } */
/* { dg-output ".*\\\[24\] OMP_DYNAMIC = 'TRUE'.*" { target native } } */
/* { dg-output ".*\\\[host] OMP_NUM_THREADS = '4,3,2'.*" { target native } } */
/* { dg-output ".*\\\[all\] OMP_NUM_THREADS = '45,46,47'.*" { target native } } */
/* { dg-output ".*\\\[device\] OMP_NUM_THREADS = '42,43,44'.*" { target native } } */
/* { dg-output ".*\\\[24\] OMP_NUM_THREADS = '14,13,12'.*" { target native } } */
/* { dg-output ".*\\\[host] OMP_SCHEDULE = 'GUIDED,2'.*" { target native } } */
/* { dg-output ".*\\\[all\] OMP_SCHEDULE = 'DYNAMIC'.*" { target native } } */
/* { dg-output ".*\\\[device\] OMP_SCHEDULE = 'GUIDED'.*" { target native } } */
/* { dg-output ".*\\\[24\] OMP_SCHEDULE = 'GUIDED,4'.*" { target native } } */
/* { dg-output ".*\\\[host] OMP_PROC_BIND = 'SPREAD'.*" { target native } } */
/* { dg-output ".*\\\[all\] OMP_PROC_BIND = 'CLOSE'.*" { target native } } */
/* { dg-output ".*\\\[device\] OMP_PROC_BIND = 'SPREAD,SPREAD'.*" { target native } } */
/* { dg-output ".*\\\[24\] OMP_PROC_BIND = 'SPREAD,CLOSE'.*" { target native } } */
/* { dg-output ".*\\\[host] OMP_STACKSIZE = '43008'.*" { target native } } */
/* { dg-output ".*\\\[all\] OMP_STACKSIZE = '44040192'.*" { target native } } */
/* { dg-output ".*\\\[device\] OMP_STACKSIZE = '44032'.*" { target native } } */
/* { dg-output ".*\\\[24\] OMP_STACKSIZE = '45056'.*" { target native } } */
/* { dg-output ".*\\\[host] OMP_WAIT_POLICY = 'ACTIVE'.*" { target native } } */
/* { dg-output ".*\\\[all\] OMP_WAIT_POLICY = 'ACTIVE'.*" { target native } } */
/* { dg-output ".*\\\[device\] OMP_WAIT_POLICY = 'PASSIVE'.*" { target native } } */
/* { dg-output ".*\\\[24\] OMP_WAIT_POLICY = 'PASSIVE'.*" { target native } } */
/* { dg-output ".*\\\[host] OMP_THREAD_LIMIT = '45'.*" { target native } } */
/* { dg-output ".*\\\[all\] OMP_THREAD_LIMIT = '43'.*" { target native } } */
/* { dg-output ".*\\\[device\] OMP_THREAD_LIMIT = '44'.*" { target native } } */
/* { dg-output ".*\\\[24\] OMP_THREAD_LIMIT = '42'.*" { target native } } */
/* { dg-output ".*\\\[host] OMP_MAX_ACTIVE_LEVELS = '42'.*" { target native } } */
/* { dg-output ".*\\\[all\] OMP_MAX_ACTIVE_LEVELS = '43'.*" { target native } } */
/* { dg-output ".*\\\[device\] OMP_MAX_ACTIVE_LEVELS = '44'.*" { target native } } */
/* { dg-output ".*\\\[24\] OMP_MAX_ACTIVE_LEVELS = '45'.*" { target native } } */
/* { dg-output ".*\\\[host] OMP_NUM_TEAMS = '42'.*" { target native } } */
/* { dg-output ".*\\\[all\] OMP_NUM_TEAMS = '43'.*" { target native } } */
/* { dg-output ".*\\\[device\] OMP_NUM_TEAMS = '44'.*" { target native } } */
/* { dg-output ".*\\\[24\] OMP_NUM_TEAMS = '45'.*" { target native } } */
/* { dg-output ".*\\\[host] OMP_TEAMS_THREAD_LIMIT = '42'.*" { target native } } */
/* { dg-output ".*\\\[all\] OMP_TEAMS_THREAD_LIMIT = '43'.*" { target native } } */
/* { dg-output ".*\\\[device\] OMP_TEAMS_THREAD_LIMIT = '44'.*" { target native } } */
/* { dg-output ".*\\\[24\] OMP_TEAMS_THREAD_LIMIT = '45'.*" { target native } } */
/* { dg-output ".*\\\[all] OMP_CANCELLATION = 'TRUE'.*" { target native } } */
/* { dg-output ".*\\\[all] OMP_DEFAULT_DEVICE = '42'.*" { target native } } */
/* { dg-output ".*\\\[all] OMP_MAX_TASK_PRIORITY = '20'.*" { target native } } */
/* { dg-output ".*\\\[all] OMP_DISPLAY_AFFINITY = 'TRUE'.*" { target native } } */
/* { dg-output ".*\\\[host] OMP_ALLOCATOR = 'omp_const_mem_alloc'.*" { target native } } */
/* { dg-output ".*\\\[all] OMP_TARGET_OFFLOAD = 'MANDATORY'.*" { target native } } */

22
libgomp/testsuite/libgomp.c-c++-common/omp-display-env-2.c Normal file
View File

@ -0,0 +1,22 @@
/* { dg-do run } */
/* { dg-set-target-env-var OMP_NUM_TEAMS "42" } */
/* This test checks if omp_display_env outputs the initial ICV values although
the value was updated. */
#include <omp.h>
#include <stdlib.h>
int
main ()
{
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'" { target native } } */