diff mbox series

[OpenACC,'exit,data'] Strip 'GOMP_MAP_STRUCT' mappings (was: [PATCH 07/13] OpenACC 2.6 deep copy: libgomp parts)

Message ID 87v9k5qxjt.fsf@euler.schwinge.homeip.net
State New
Headers show
Series [OpenACC,'exit,data'] Strip 'GOMP_MAP_STRUCT' mappings (was: [PATCH 07/13] OpenACC 2.6 deep copy: libgomp parts) | expand

Commit Message

Thomas Schwinge June 5, 2020, 4:36 p.m. UTC 
Hi!

On 2020-05-20T11:37:35+0200, I wrote:
> 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 <julian@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 [...]

> Further citing myself,
> <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: [...]

> 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, [...]
>
> 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': [...]

> ..., 'deep-copy-7.c.005t.gimple' gets simplified as expected: [...]

> [...] 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...

You've picked that up in your
<http://mid.mail-archive.com/cf6237ce3f7cad564e2581e450c7d0346f53365f.1591276990.git.julian@codesourcery.com>
patch submission.  Thanks for the test case you've included, I changed
that one just a bit: '-fdump-tree-gimple' scanning (as that's where the
processing happens), removed the 'stdlib.h' dependency, made the 'struct'
(and thus the 'GOMP_MAP_STRUCT' that's generated) a bit more interesting.
I've pushed "[OpenACC 'exit data'] Strip 'GOMP_MAP_STRUCT' mappings" to
master branch in commit 1afc4672561a41dfbf4e3f2c1f35f7a5b7a20339, and
releases/gcc-10 branch in commit
27e985479e931771472663cad34f8b99c6f62627, see attached.


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
diff mbox series

Patch

From 27e985479e931771472663cad34f8b99c6f62627 Mon Sep 17 00:00:00 2001
From: Thomas Schwinge <thomas@codesourcery.com>
Date: Wed, 20 May 2020 10:56:55 +0200
Subject: [PATCH] [OpenACC 'exit data'] Strip 'GOMP_MAP_STRUCT' mappings

These are not itself necessary for OpenACC 'exit data' directives, and are
skipped over (now) in libgomp.  We might as well not emit them to start with,
in line with the equivalent OpenMP directive.  We keep the no-op handling in
libgomp for the reason of backward compatibility.

	gcc/
	* gimplify.c (gimplify_adjust_omp_clauses): Remove
	'GOMP_MAP_STRUCT' mapping from OpenACC 'exit data' directives.
	gcc/testsuite/
	* c-c++-common/goacc/struct-enter-exit-data-1.c: New file.
	libgomp/
	* oacc-mem.c (goacc_exit_data_internal) <GOMP_MAP_STRUCT>: Explain
	special handling.

Co-Authored-By: Julian Brown <julian@codesourcery.com>
(cherry picked from commit 1afc4672561a41dfbf4e3f2c1f35f7a5b7a20339)
---
 gcc/gimplify.c                                |  3 ++-
 .../goacc/struct-enter-exit-data-1.c          | 27 +++++++++++++++++++
 libgomp/oacc-mem.c                            |  5 ++--
 3 files changed, 32 insertions(+), 3 deletions(-)
 create mode 100644 gcc/testsuite/c-c++-common/goacc/struct-enter-exit-data-1.c

diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index 0bac99002102..4b3306cfc258 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -10408,7 +10408,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
diff --git a/gcc/testsuite/c-c++-common/goacc/struct-enter-exit-data-1.c b/gcc/testsuite/c-c++-common/goacc/struct-enter-exit-data-1.c
new file mode 100644
index 000000000000..df405e448b28
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/struct-enter-exit-data-1.c
@@ -0,0 +1,27 @@ 
+/* Check 'GOMP_MAP_STRUCT' mapping, and in particular that it gets removed from
+   OpenACC 'exit data' directives.  */
+
+/* { dg-additional-options "-fdump-tree-gimple" } */
+
+struct str {
+  int a;
+  int *b;
+  int *c;
+  int d;
+  int *e;
+  int f;
+};
+
+#define N 1024
+
+void
+test (int *b, int *c, int *e)
+{
+  struct str s = { .a = 0, .b = b, .c = c, .d = 0, .e = e, .f = 0 };
+
+#pragma acc enter data copyin(s.a, s.b[0:N], s.c[0:N] /* , s.d */ /* , s.e[0:N] */, s.f)
+  /* { dg-final { scan-tree-dump {(?n)#pragma omp target oacc_enter_exit_data map\(struct:s \[len: 4\]\) map\(to:s.a \[len: [0-9]+\]\) map\(alloc:s.b \[len: [0-9]+\]\) map\(alloc:s.c \[len: [0-9]+\]\) map\(to:s.f \[len: [0-9]+\]\) map\(to:\*[_0-9]+ \[len: [0-9]+\]\) map\(attach:s.b \[bias: 0\]\) map\(to:\*[_0-9]+ \[len: [0-9]+\]\) map\(attach:s.c \[bias: 0\]\)$} gimple } } */
+
+#pragma acc exit data copyout(s.a, s.b[0:N], s.c[0:N] /* , s.d */ /* , s.e[0:N] */, s.f)
+  /* { dg-final { scan-tree-dump {(?n)#pragma omp target oacc_enter_exit_data map\(from:s.a \[len: [0-9]+\]\) map\(release:s.b \[len: [0-9]+\]\) map\(release:s.c \[len: [0-9]+\]\) map\(from:s.f \[len: [0-9]+\]\) map\(from:\*[_0-9]+ \[len: [0-9]+\]\) map\(detach:s.b \[bias: 0\]\) map\(from:\*[_0-9]+ \[len: [0-9]+\]\) map\(detach:s.c \[bias: 0\]\)$} gimple } } */
+}
diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c
index 1e3685a073da..936ae649dd93 100644
--- a/libgomp/oacc-mem.c
+++ b/libgomp/oacc-mem.c
@@ -1181,8 +1181,9 @@  goacc_exit_data_internal (struct gomp_device_descr *acc_dev, size_t mapnum,
 
 	case GOMP_MAP_STRUCT:
 	  /* Skip the 'GOMP_MAP_STRUCT' itself, and use the regular processing
-	     for all its entries.  TODO: don't generate these no-op
-	     'GOMP_MAP_STRUCT's.  */
+	     for all its entries.  This special handling exists for GCC 10.1
+	     compatibility; afterwards, we're not generating these no-op
+	     'GOMP_MAP_STRUCT's anymore.  */
 	  break;
 
 	default:
-- 
2.26.2