Hi Julian!
Moving this over, from the "Fix component mappings with derived types for
OpenACC" thread,
<http://mid.mail-archive.com/20200110014945.5643ace5@squid.athome>, where
you propose to change this 'GOMP_MAP_STRUCT' handling code:
On 2019-12-17T22:03:47-0800, Julian Brown <jul...@codesourcery.com> wrote:
> --- a/libgomp/oacc-mem.c
> +++ b/libgomp/oacc-mem.c
> @@ -1075,6 +1119,39 @@ goacc_exit_data_internal (struct gomp_device_descr
> *acc_dev, size_t mapnum,
> gomp_remove_var_async (acc_dev, n, aq);
> }
> break;
> +
> + case GOMP_MAP_STRUCT:
> + {
> + int elems = sizes[i];
> + for (int j = 1; j <= elems; j++)
> + {
> + struct splay_tree_key_s k;
> + k.host_start = (uintptr_t) hostaddrs[i + j];
> + k.host_end = k.host_start + sizes[i + j];
> + splay_tree_key str;
> + str = splay_tree_lookup (&acc_dev->mem_map, &k);
> + if (str)
> + {
> + if (finalize)
> + {
> + str->refcount -= str->virtual_refcount;
> + str->virtual_refcount = 0;
> + }
> + if (str->virtual_refcount > 0)
> + {
> + str->refcount--;
> + str->virtual_refcount--;
> + }
> + else if (str->refcount > 0)
> + str->refcount--;
> + if (str->refcount == 0)
> + gomp_remove_var_async (acc_dev, str, aq);
> + }
> + }
> + i += elems;
> + }
> + break;
> +
> default:
> gomp_fatal (">>>> goacc_exit_data_internal UNHANDLED kind 0x%.2x",
> kind);
... into an "empty 'case GOMP_MAP_STRUCT:' as a no-op, so that we don't
run into 'default:' case 'goacc_exit_data_internal UNHANDLED kind'" (my
words/interpretation).
Further citing myself,
<87ftbw9kqh.fsf@euler.schwinge.homeip.net">http://mid.mail-archive.com/87ftbw9kqh.fsf@euler.schwinge.homeip.net>:
| Is my understanding correct that "fixed" GCC won't generate such
| 'GOMP_MAP_STRUCT' anymore (I have't studied in detail), and this empty
| 'case GOMP_MAP_STRUCT:' only remains in here for backwards compatibility?
| In this case, please add a comment to the code, stating this.
My guess was wrong: running a quick test, I do see that we still generate
'GOMP_MAP_STRUCT' for OpenACC unmap:
--- libgomp/oacc-mem.c
+++ libgomp/oacc-mem.c
@@ -1163,6 +1165,7 @@ goacc_exit_data_internal
break;
case GOMP_MAP_STRUCT:
+ assert (!"GOMP_MAP_STRUCT");
break;
default:
gomp_fatal (">>>> goacc_exit_data_internal UNHANDLED kind 0x%.2x",
kind);
... regresses:
- 'libgomp.oacc-c-c++-common/deep-copy-7.c'
- 'libgomp.oacc-c-c++-common/deep-copy-8.c'
- 'libgomp.oacc-fortran/classtypes-2.f95'
- 'libgomp.oacc-fortran/deep-copy-4.f90'
- 'libgomp.oacc-fortran/deep-copy-5.f90'
- 'libgomp.oacc-fortran/deep-copy-6.f90'
- 'libgomp.oacc-fortran/derivedtype-2.f95'
| Otherwise,
| please add a comment why "do nothing" is appropriate for
| 'GOMP_MAP_STRUCT'.
I suppose we still need to unmap the "'GOMP_MAP_STRUCT' components", but
can do that individually, outside of the 'GOMP_MAP_STRUCT' context.
That'd then also explain...
| In particular, for both scenarios, why we don't need
| to skip the following 'sizes[i]' mappings?
... this question.
But one step back. Why generate 'GOMP_MAP_STRUCT' for unmap, if we then
just skip it in libgomp handling? Cross checking, OpenMP
'libgomp/target.c:gomp_exit_data' also doesn't expect to see any
'GOMP_MAP_STRUCT'.
For example, 'libgomp.oacc-c-c++-common/deep-copy-7.c':
#pragma acc exit data copyout(v.b[:n]) finalize
#pragma acc exit data delete(v.a)
'deep-copy-7.c.004t.original':
#pragma acc exit data finalize map(from:*v.b [len: (sizetype) n * 4])
map(attach_detach:v.b [bias: 0]);
#pragma acc exit data map(release:v.a);
'deep-copy-7.c.005t.gimple':
#pragma omp target oacc_enter_exit_data finalize map(struct:v [len: 1])
map(delete:v.b [len: 8]) map(force_from:*_15 [len: _14]) map(force_detach:v.b
[bias: 0])
#pragma omp target oacc_enter_exit_data map(struct:v [len: 1])
map(release:v.a [len: 4])
I haven't studied 'GOMP_MAP_STRUCT' in detail (so it may be easy to prove
me wrong), but for OpenACC 'exit data' etc., these 'map(struct:[...])'
seem "pointless" to me, given that in libgomp we (intend to) just skip
them?
Well, and quickly I find that's exactly what OpenMP 'target exit data' is
doing, and doing the same for OpenACC 'exit data':
--- gcc/gimplify.c
+++ gcc/gimplify.c
@@ -10406,7 +10406,8 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p,
gimple_seq body, tree *list_p,
}
}
else if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_STRUCT
- && code == OMP_TARGET_EXIT_DATA)
+ && (code == OMP_TARGET_EXIT_DATA
+ || code == OACC_EXIT_DATA))
remove = true;
else if (DECL_SIZE (decl)
&& TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST
..., 'deep-copy-7.c.005t.gimple' gets simplified as expected:
- #pragma omp target oacc_enter_exit_data finalize map(struct:v
[len: 1]) map(delete:v.b [len: 8]) map(force_from:*_15 [len: _14])
map(force_detach:v.b [bias: 0])
+ #pragma omp target oacc_enter_exit_data finalize map(delete:v.b
[len: 8]) map(force_from:*_15 [len: _14]) map(force_detach:v.b [bias: 0])
- #pragma omp target oacc_enter_exit_data map(struct:v [len: 1])
map(release:v.a [len: 4])
+ #pragma omp target oacc_enter_exit_data map(release:v.a [len: 4])
..., and all the "'assert' regressions" mentioned above again disappear,
and so 'GOMP_MAP_STRUCT' handling could be removed from
'libgomp/oacc-mem.c:goacc_exit_data_internal' (and 'default:'
'gomp_fatal' would then catch any such cases).
But of course, given that GCC 10.1 now does generate these
'GOMP_MAP_STRUCT's, we do have to support them in one way or another...
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
