Hi!
On 2019-12-17T22:03:47-0800, Julian Brown <jul...@codesourcery.com> wrote:
> --- a/libgomp/oacc-mem.c
> +++ b/libgomp/oacc-mem.c
> static int
> -find_group_last (int pos, size_t mapnum, unsigned short *kinds)
> +find_group_last (int pos, size_t mapnum, size_t *sizes, unsigned short
> *kinds)
> {
> unsigned char kind0 = kinds[pos] & 0xff;
> - int first_pos = pos, last_pos = pos;
> + int first_pos = pos;
>
> - if (kind0 == GOMP_MAP_TO_PSET)
> + switch (kind0)
> {
> + case GOMP_MAP_TO_PSET:
> while (pos + 1 < mapnum && (kinds[pos + 1] & 0xff) == GOMP_MAP_POINTER)
> - last_pos = ++pos;
> + pos++;
> /* We expect at least one GOMP_MAP_POINTER after a GOMP_MAP_TO_PSET.
> */
> - assert (last_pos > first_pos);
> - }
> - else
> - {
> + assert (pos > first_pos);
> + break;
> +
> + case GOMP_MAP_STRUCT:
> + pos += sizes[pos];
> + break;
> +
> + case GOMP_MAP_POINTER:
> + case GOMP_MAP_ALWAYS_POINTER:
> + /* These mappings are only expected after some other mapping. If we
> + see one by itself, something has gone wrong. */
> + gomp_fatal ("unexpected mapping");
> + break;
> +
> + default:
> /* GOMP_MAP_ALWAYS_POINTER can only appear directly after some other
> mapping. */
> - if (pos + 1 < mapnum
> - && (kinds[pos + 1] & 0xff) == GOMP_MAP_ALWAYS_POINTER)
> - return pos + 1;
> + if (pos + 1 < mapnum)
> + {
> + unsigned char kind1 = kinds[pos + 1] & 0xff;
> + if (kind1 == GOMP_MAP_ALWAYS_POINTER)
> + return pos + 1;
> + }
>
> - /* We can have one or several GOMP_MAP_POINTER mappings after a to/from
> + /* We can have zero or more GOMP_MAP_POINTER mappings after a to/from
> (etc.) mapping. */
> while (pos + 1 < mapnum && (kinds[pos + 1] & 0xff) == GOMP_MAP_POINTER)
> - last_pos = ++pos;
> + pos++;
> }
>
> - return last_pos;
> + return pos;
> }
So this now causes grouped (!) mapping of all of 'GOMP_MAP_STRUCT', that
is, all its "members" at once.
This, I suppose, mandated the removal of (some of) the 'is_tgt_unmapped'
checking (unfortunately committed not here, but as part of r279621
"OpenACC reference count overhaul"), where we had unmapping code
(conceptually) similar to:
bool is_tgt_unmapped = gomp_remove_var (acc_dev, n);
assert (is_tgt_unmapped);
I'd introduced this a little bit earlier, finding this a simple yet
effective run-time, low-overhead consistency checking of (certain aspects
of) reference counting -- so just noting here that it's somewhat bad that
we can't have this anymore "just" because of 'GOMP_MAP_STRUCT'. (Maybe
there is a way to get it back; that's for later?)
Anyway, the code changes were incomplete, consider:
#include <assert.h>
#include <openacc.h>
struct s
{
char a;
char b;
};
int main ()
{
struct s s;
#pragma acc enter data create(s.a, s.b)
assert (acc_is_present (&s.a, sizeof s.a));
assert (acc_is_present (&s.b, sizeof s.b));
#if 0
// works
# pragma acc exit data delete(s.a)
# pragma acc exit data delete(s.b)
#else
acc_delete (&s.a, sizeof s.a); // fails
acc_delete (&s.b, sizeof s.b);
#endif
assert (!acc_is_present (&s.a, sizeof s.a));
assert (!acc_is_present (&s.b, sizeof s.b));
return 0;
}
The 'acc_delete' variant exercises a code path that still contains the
'is_tgt_unmapped' checking, and that triggers for 'acc_delete ([s.a])'
then, as 's.a' has been mapped in one group together with 's.b', which
remains mapped until 'acc_delete ([s.b])', and thus
"[...]/libgomp/oacc-mem.c:726: goacc_exit_datum: Assertion
`is_tgt_unmapped' failed". Unless anybody has a better plan for this
whole topic, I'll commit the obvious (removal), later (next week).
--- libgomp/oacc-mem.c
+++ libgomp/oacc-mem.c
@@ -724,7 +724,11 @@ goacc_exit_datum (void *h, size_t s, unsigned short
kind, int async)
else
{
bool is_tgt_unmapped = gomp_remove_var (acc_dev, n);
+#if 0
assert (is_tgt_unmapped);
+#else
+ (void) is_tgt_unmapped;
+#endif
}
}
Grüße
Thomas
-----------------
Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany
Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander
Walter
