On Fri, Oct 18, 2019 at 11:27:39AM +0200, Tobias Burnus wrote:
> Currently, one has for
> !$omp target exit data map(delete:x)
> in the original dump:
> #pragma omp target exit data map(delete:*x) map(alloc:x [pointer assign,
> bias: 0])
>
> The "alloc:" not only does not make sense but also gives run-time messages
> like:
> libgomp: GOMP_target_enter_exit_data unhandled kind 0x04
>
> [Depending on the data type, in gfc_trans_omp_clauses's OMP_LIST_MAP, add
> map clauses of type GOMP_MAP_POINTER and/or GOMP_MAP_TO_PSET.]
>
> That's for release:/delete:. However, for 'target exit data'
> (GOMP_target_enter_exit_data) the same issue occurs for "from:"/"always,
> from:". But "from:" implies "alloc:". – While "alloc:" does not make sense
> for "target exit data" or "update", for "target" or "target data" it surely
> matters. Hence, I only exclude "from:" for exit data and update.
>
> See attached patch. I have additionally Fortran-fied libgomp.c/target-20.c
> to have at least one 'enter/exit target data' test case for Fortran.
>
> Build + regtested on x86_64-gnu-linux w/o offloading. And I have tested the
> new test case with nvptx.
I believe it is easier to handle it at the same spot as we do it e.g. for
C/C++ pointer attachments (where we create the same clauses regardless of
the exact construct and then drop them later), in particular in
gimplify_scan_omp_clauses.
There we have:
case OMP_TARGET:
break;
case OACC_DATA:
if (TREE_CODE (TREE_TYPE (decl)) != ARRAY_TYPE)
break;
/* FALLTHRU */
case OMP_TARGET_DATA:
case OMP_TARGET_ENTER_DATA:
case OMP_TARGET_EXIT_DATA:
case OACC_ENTER_DATA:
case OACC_EXIT_DATA:
case OACC_HOST_DATA:
if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER
|| (OMP_CLAUSE_MAP_KIND (c)
== GOMP_MAP_FIRSTPRIVATE_REFERENCE))
/* For target {,enter ,exit }data only the array slice is
mapped, but not the pointer to it. */
remove = true;
break;
So, I think best would be to add
if (code == OMP_TARGET_EXIT_DATA && OMP_CLAUSE_MAP_KIND (c) ==
GOMP_MAP_WHATEVER_IS_NOT_VALID_FOR_EXIT_DATA) remove = true;
with a comment explaining that.
The testcase LGTM.
Jakub
