On Sun, Oct 09, 2022 at 02:51:37PM -0700, Julian Brown wrote:
> This patch adds support for non-constant component offsets in "map"
> clauses for OpenMP (and the equivalants for OpenACC), which are not able
> to be sorted into order at compile time. Normally struct accesses in
> such clauses are gathered together and sorted into increasing address
> order after a "GOMP_MAP_STRUCT" node: if we have variable indices,
> that is no longer possible.
>
> This patch adds support for such mappings by adding a new variant of
> GOMP_MAP_STRUCT that does not require the list of nodes following to
> be in sorted order. This passes right down to the runtime: the list is
> sorted in libgomp according to the dynamic values of the offsets after
> the newly-introduced GOMP_MAP_STRUCT_UNORD node.
>
> This mostly affects arrays of structs indexed by variables in C and C++,
> but can also affect derived-type arrays with constant indexes when they
> have an array descriptor.
I don't understand why this is needed.
We have a restriction that ought to make all such cases invalid:
"If multiple list items are explicitly mapped on the same construct and have
the same containing
array or have base pointers that share original storage, and if any of the list
items do not have
corresponding list items that are present in the device data environment prior
to a task
encountering the construct, then the list items must refer to the same array
elements of either the
containing array or the implicit array of the base pointers."
So, all those
#pragma omp target map(t->a[i].p, t->a[j].p) etc. cases are invalid unless
i == j, so IMNSHO one doesn't need to worry about unordered cases.
> +#if defined(_GNU_SOURCE) || defined(__GNUC__)
> +static int
> +compare_addr_r (const void *a, const void *b, void *data)
> +{
> + void **hostaddrs = (void **) data;
> + int ai = *(int *) a, bi = *(int *) b;
> + if (hostaddrs[ai] < hostaddrs[bi])
> + return -1;
> + else if (hostaddrs[ai] > hostaddrs[bi])
> + return 1;
> + return 0;
> +}
> +#endif
Note, not all glibcs have qsort_r and _GNU_SOURCE macro being defined
doesn't imply glibc. You'd need something like _GLIBC_PREREQ macro
and require 2.8 or later.
> +
> static inline __attribute__((always_inline)) struct target_mem_desc *
> gomp_map_vars_internal (struct gomp_device_descr *devicep,
> struct goacc_asyncqueue *aq, size_t mapnum,
> @@ -968,6 +982,17 @@ gomp_map_vars_internal (struct gomp_device_descr
> *devicep,
> tgt->device_descr = devicep;
> tgt->prev = NULL;
> struct gomp_coalesce_buf cbuf, *cbufp = NULL;
> + size_t hostaddr_idx;
> +
> +#if !defined(_GNU_SOURCE) && defined(__GNUC__)
> + /* If we don't have _GNU_SOURCE (thus no qsort_r), but we are compiling
> with
> + GCC (and why wouldn't we be?), we can use this nested function for
> + regular qsort. */
> + int compare_addr (const void *a, const void *b)
> + {
> + return compare_addr_r (a, b, (void *) &hostaddrs[hostaddr_idx]);
> + }
> +#endif
Please don't use nested functions in libgomp.
> + int *order = NULL;
> + if ((kind & typemask) == GOMP_MAP_STRUCT_UNORD)
> + {
> + order = (int *) gomp_alloca (sizeof (int) * sizes[i]);
> + for (int j = 0; j < sizes[i]; j++)
> + order[j] = j;
> +#ifdef _GNU_SOURCE
> + qsort_r (order, sizes[i], sizeof (int), &compare_addr_r,
> + &hostaddrs[i + 1]);
> +#elif defined(__GNUC__)
> + hostaddr_idx = i + 1;
> + qsort (order, sizes[i], sizeof (int), &compare_addr);
> +#else
> +#error no threadsafe qsort
> +#endif
This is too ugly. If this is really needed (see above) and
you need fallback for missing qsort_t, better sort array of tuples
containing the order number and some data pointer needed for the comparison
routine.
Jakub
