Re: [Patch][Fortran/OpenMP] Don't create "alloc:" for 'target exit data'

classic Classic list List threaded Threaded
3 messages Options
Reply | Threaded
Open this post in threaded view
|

Re: [Patch][Fortran/OpenMP] Don't create "alloc:" for 'target exit data'

Jakub Jelinek
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

Reply | Threaded
Open this post in threaded view
|

Re: [Patch][Fortran/OpenMP] Don't create "alloc:" for 'target exit data'

Tobias Burnus-3
On 10/30/19 11:12 AM, Jakub Jelinek wrote:
> 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. […]

I concur. Semantically, it is not identical – but I think still okay.

For 'omp exit data', 'to:'/'alloc:' mapping does not make sense and it
not handled in libgomp's gomp_exit_data. Hence, I exclude
GOMP_MAP_POINTER (dump: 'alloc:') and GOMP_MAP_TO_PSET (dump: 'to:'). –
Those are only internally used, hence, user-specified 'alloc:' will get
diagnosed.

['delete:'/'release:' in other directives than 'exit data' doesn't make
much sense. Other directives accept it but their libgomp function
silently ignore it.]

'omp update': The gomp_update function only handles GOMP_MAP_COPY_TO_P
and GOMP_MAP_COPY_FROM_P (and silently ignores others). Both macros have
!((X) & GOMP_MAP_FLAG_SPECIAL). Hence, we can save a few bytes and avoid
calling 'omp update' with GOMP_MAP_POINTER and GOMP_MAP_TO_PSET.

[TO_PSET only appears in gfc_trans_omp_clauses (once); POINTER appears
there and in gfc_omp_finish_clause and in c/c-typeck.c's
handle_omp_array_sections but only if "(ort != C_ORT_OMP && ort !=
C_ORT_ACC)".]


I moved trans-openmp.c change to gimplify.c and left the test case
unchanged. Then, I bootstrapped on a non-offloading system and regtested
it also with a nvptx system.

Tobias


exit-target-data-v3.diff (4K) Download Attachment
Reply | Threaded
Open this post in threaded view
|

Re: [Patch][Fortran/OpenMP] Don't create "alloc:" for 'target exit data'

Thomas Schwinge-8
Hi!

On 2019-10-30T16:48:43+0100, Tobias Burnus <[hidden email]> wrote:
> --- /dev/null
> +++ b/libgomp/testsuite/libgomp.fortran/target9.f90

As obvious; see attached, committed "Torture testing:
'libgomp.fortran/target9.f90'" to trunk in r278045.


Grüße
 Thomas



From d462cbc6c489949752b4d652abec30dbb95c2855 Mon Sep 17 00:00:00 2001
From: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
Date: Mon, 11 Nov 2019 08:50:40 +0000
Subject: [PATCH] Torture testing: 'libgomp.fortran/target9.f90'

        libgomp/
        * testsuite/libgomp.fortran/target9.f90: Specify 'dg-do run'.

git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/trunk@278045 138bc75d-0d04-0410-961f-82ee72b054a4
---
 libgomp/ChangeLog                             | 2 ++
 libgomp/testsuite/libgomp.fortran/target9.f90 | 1 +
 2 files changed, 3 insertions(+)

diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog
index 0e73cadb6cd0..1fc8c471b6f8 100644
--- a/libgomp/ChangeLog
+++ b/libgomp/ChangeLog
@@ -1,5 +1,7 @@
 2019-11-11  Thomas Schwinge  <[hidden email]>
 
+ * testsuite/libgomp.fortran/target9.f90: Specify 'dg-do run'.
+
  * testsuite/libgomp.fortran/use_device_addr-3.f90: Specify 'dg-do
  run'.
  * testsuite/libgomp.fortran/use_device_addr-4.f90: Likewise.
diff --git a/libgomp/testsuite/libgomp.fortran/target9.f90 b/libgomp/testsuite/libgomp.fortran/target9.f90
index 91d60a33307e..30adc1bd70af 100644
--- a/libgomp/testsuite/libgomp.fortran/target9.f90
+++ b/libgomp/testsuite/libgomp.fortran/target9.f90
@@ -1,3 +1,4 @@
+! { dg-do run }
 ! { dg-require-effective-target offload_device_nonshared_as } */
 
 module target_test
--
2.17.1


signature.asc (671 bytes) Download Attachment