Make possible 'scan-tree-dump' of 'lower_omp_target' mapping kinds

Message ID 8736l1qi3o.fsf@euler.schwinge.homeip.net
State New
Headers show
Series
  • Make possible 'scan-tree-dump' of 'lower_omp_target' mapping kinds
Related show

Commit Message

Thomas Schwinge May 26, 2019, 4:46 p.m.
Hi!

To establish some suitable testsuite coverage for a task that I'm working
on, I need to do 'scan-tree-dump' of 'lower_omp_target' mapping kinds.
Is the attached OK?

Any suggestions about whether/how to restrict the (effective?) targets
this gets run for, because no doubt there are target-specific bits at
least in the alignment chosen.  The attached test case passes for
x86_64-pc-linux-gnu with '--target_board=unix\{,-m32,-mx32\}'.  (I didn't
verify the mappings generated, but just documented the status quo.)

If approving this patch, please respond with "Reviewed-by: NAME <EMAIL>"
so that your effort will be recorded in the commit log, see
<https://gcc.gnu.org/wiki/Reviewed-by>.


Grüße
 Thomas

Comments

Jakub Jelinek May 27, 2019, 4:51 p.m. | #1
On Sun, May 26, 2019 at 06:46:19PM +0200, Thomas Schwinge wrote:
> To establish some suitable testsuite coverage for a task that I'm working

> on, I need to do 'scan-tree-dump' of 'lower_omp_target' mapping kinds.

> Is the attached OK?

> 

> Any suggestions about whether/how to restrict the (effective?) targets

> this gets run for, because no doubt there are target-specific bits at

> least in the alignment chosen.  The attached test case passes for

> x86_64-pc-linux-gnu with '--target_board=unix\{,-m32,-mx32\}'.  (I didn't

> verify the mappings generated, but just documented the status quo.)


The arrays are emitted in the *.omplower dump, so I think it is much better
to scan-tree-dump their content if for whatever reason adding a runtime
testcase isn't sufficient over adding further printouts and matching that.

	Jakub
Thomas Schwinge May 27, 2019, 7:05 p.m. | #2
Hi!

On Mon, 27 May 2019 18:51:22 +0200, Jakub Jelinek <jakub@redhat.com> wrote:
> On Sun, May 26, 2019 at 06:46:19PM +0200, Thomas Schwinge wrote:

> > To establish some suitable testsuite coverage for a task that I'm working

> > on, I need to do 'scan-tree-dump' of 'lower_omp_target' mapping kinds.

> > Is the attached OK?

> > 

> > Any suggestions about whether/how to restrict the (effective?) targets

> > this gets run for, because no doubt there are target-specific bits at

> > least in the alignment chosen.  The attached test case passes for

> > x86_64-pc-linux-gnu with '--target_board=unix\{,-m32,-mx32\}'.  (I didn't

> > verify the mappings generated, but just documented the status quo.)

> 

> The arrays are emitted in the *.omplower dump, so I think it is much better

> to scan-tree-dump their content


That's not feasible in the general case.

> if for whatever reason adding a runtime

> testcase isn't sufficient


Why should I have an execution test for something that should really be
verified at the compiler side, at compile time?  And, for example, an
execution test cannot detect if a 'GOMP_MAP_FIRSTPRIVATE_INT' is being
used (instead of something else).

> over adding further printouts and matching that.


I had assumed that you'd noticed that a lot of compiler passes are
dumping stuff that can then be scanned for.  Are you insisting that GCC's
OMP code must be complex and unmaintainable?


Current version of that patch is attached: changed to generally print
'kind' first, and don't print non-existing 'align' for
'OMP_CLAUSE_USE_DEVICE_PTR', 'OMP_CLAUSE_IS_DEVICE_PTR'.

If approving this patch, please respond with "Reviewed-by: NAME <EMAIL>"
so that your effort will be recorded in the commit log, see
<https://gcc.gnu.org/wiki/Reviewed-by>.


Grüße
 Thomas
From 30a64c6cf7b8096f27836c62f434eb4bc0ebe966 Mon Sep 17 00:00:00 2001
From: Thomas Schwinge <thomas@codesourcery.com>

Date: Mon, 27 May 2019 18:21:15 +0200
Subject: [PATCH] Make possible 'scan-tree-dump' of 'lower_omp_target' mapping
 kinds

	gcc/
	* omp-low.c (lower_omp_target): Dump mapping kinds.
	gcc/testsuite/
	* c-c++-common/gomp/lower_omp_target-mappings-1.c: New file.
---
 gcc/omp-low.c                                 | 102 +++++++++++++++---
 .../gomp/lower_omp_target-mappings-1.c        |  50 +++++++++
 2 files changed, 136 insertions(+), 16 deletions(-)
 create mode 100644 gcc/testsuite/c-c++-common/gomp/lower_omp_target-mappings-1.c

diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index faab5d384280..bbeac7f61846 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -9218,15 +9218,23 @@ lower_omp_taskreg (gimple_stmt_iterator *gsi_p, omp_context *ctx)
 static void
 lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
 {
+  pretty_printer pp;
   tree clauses;
   tree child_fn, t, c;
   gomp_target *stmt = as_a <gomp_target *> (gsi_stmt (*gsi_p));
   gbind *tgt_bind, *bind, *dep_bind = NULL;
   gimple_seq tgt_body, olist, ilist, fplist, new_body;
   location_t loc = gimple_location (stmt);
+  const char *loc_str = NULL;
   bool offloaded, data_region;
   unsigned int map_cnt = 0;
 
+  if (dump_file && (dump_flags & TDF_DETAILS))
+    {
+      dump_location (&pp, loc);
+      loc_str = pp_formatted_text (&pp);
+    }
+
   offloaded = is_gimple_omp_offloaded (stmt);
   switch (gimple_omp_target_kind (stmt))
     {
@@ -9687,7 +9695,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
 	    if (TREE_CODE (s) != INTEGER_CST)
 	      TREE_STATIC (TREE_VEC_ELT (t, 1)) = 0;
 
-	    unsigned HOST_WIDE_INT tkind, tkind_zero;
+	    unsigned HOST_WIDE_INT tkind, tkind_align;
+	    unsigned HOST_WIDE_INT tkind_zero, tkind_zero_align;
 	    switch (OMP_CLAUSE_CODE (c))
 	      {
 	      case OMP_CLAUSE_MAP:
@@ -9743,25 +9752,43 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
 				 < (HOST_WIDE_INT_C (1U) << talign_shift));
 	    gcc_checking_assert (tkind_zero
 				 < (HOST_WIDE_INT_C (1U) << talign_shift));
-	    talign = ceil_log2 (talign);
-	    tkind |= talign << talign_shift;
-	    tkind_zero |= talign << talign_shift;
-	    gcc_checking_assert (tkind
+	    {
+	      unsigned int talign2 = ceil_log2 (talign);
+	      tkind_align = tkind | (talign2 << talign_shift);
+	      tkind_zero_align = tkind_zero | (talign2 << talign_shift);
+	    }
+	    gcc_checking_assert (tkind_align
 				 <= tree_to_uhwi (TYPE_MAX_VALUE (tkind_type)));
-	    gcc_checking_assert (tkind_zero
+	    gcc_checking_assert (tkind_zero_align
 				 <= tree_to_uhwi (TYPE_MAX_VALUE (tkind_type)));
-	    if (tkind == tkind_zero)
-	      x = build_int_cstu (tkind_type, tkind);
+	    if (tkind_align == tkind_zero_align)
+	      x = build_int_cstu (tkind_type, tkind_align);
 	    else
 	      {
 		TREE_STATIC (TREE_VEC_ELT (t, 2)) = 0;
 		x = build3 (COND_EXPR, tkind_type,
 			    fold_build2 (EQ_EXPR, boolean_type_node,
 					 unshare_expr (s), size_zero_node),
-			    build_int_cstu (tkind_type, tkind_zero),
-			    build_int_cstu (tkind_type, tkind));
+			    build_int_cstu (tkind_type, tkind_zero_align),
+			    build_int_cstu (tkind_type, tkind_align));
 	      }
 	    CONSTRUCTOR_APPEND_ELT (vkind, purpose, x);
+
+	    if (dump_file && (dump_flags & TDF_DETAILS))
+	      {
+		fprintf (dump_file, "Mapping %s%s [%u] '",
+			 loc_str,
+			 IDENTIFIER_POINTER (DECL_NAME (ctx->cb.dst_fn)),
+			 (unsigned int) tree_to_uhwi (purpose));
+		print_generic_expr (dump_file, ovar, dump_flags);
+		fprintf (dump_file, "': kind = %u/%u, size = ",
+			 (unsigned int) tkind,
+			 (unsigned int) tkind_zero);
+		print_generic_expr (dump_file, s, dump_flags);
+		fprintf (dump_file, ", align = %u\n",
+			 talign);
+	      }
+
 	    if (nc && nc != c)
 	      c = nc;
 	    break;
@@ -9826,12 +9853,29 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
 
 	    gcc_checking_assert (tkind
 				 < (HOST_WIDE_INT_C (1U) << talign_shift));
-	    talign = ceil_log2 (talign);
-	    tkind |= talign << talign_shift;
-	    gcc_checking_assert (tkind
+	    {
+	      unsigned int talign2 = ceil_log2 (talign);
+	      tkind_align = tkind | (talign2 << talign_shift);
+	    }
+	    gcc_checking_assert (tkind_align
 				 <= tree_to_uhwi (TYPE_MAX_VALUE (tkind_type)));
 	    CONSTRUCTOR_APPEND_ELT (vkind, purpose,
-				    build_int_cstu (tkind_type, tkind));
+				    build_int_cstu (tkind_type, tkind_align));
+
+	    if (dump_file && (dump_flags & TDF_DETAILS))
+	      {
+		fprintf (dump_file, "Mapping %s%s [%u] '",
+			 loc_str,
+			 IDENTIFIER_POINTER (DECL_NAME (ctx->cb.dst_fn)),
+			 (unsigned int) tree_to_uhwi (purpose));
+		print_generic_expr (dump_file, ovar, dump_flags);
+		fprintf (dump_file, "': kind = %u, size = ",
+			 (unsigned int) tkind);
+		print_generic_expr (dump_file, s, dump_flags);
+		fprintf (dump_file, ", align = %u\n",
+			 talign);
+	      }
+
 	    break;
 
 	  case OMP_CLAUSE_USE_DEVICE_PTR:
@@ -9862,14 +9906,40 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
 	    CONSTRUCTOR_APPEND_ELT (vsize, purpose, s);
 	    gcc_checking_assert (tkind
 				 < (HOST_WIDE_INT_C (1U) << talign_shift));
-	    gcc_checking_assert (tkind
+	    tkind_align = tkind;
+	    gcc_checking_assert (tkind_align
 				 <= tree_to_uhwi (TYPE_MAX_VALUE (tkind_type)));
 	    CONSTRUCTOR_APPEND_ELT (vkind, purpose,
-				    build_int_cstu (tkind_type, tkind));
+				    build_int_cstu (tkind_type, tkind_align));
+
+	    if (dump_file && (dump_flags & TDF_DETAILS))
+	      {
+		fprintf (dump_file, "Mapping %s%s [%u] '",
+			 loc_str,
+			 IDENTIFIER_POINTER (DECL_NAME (ctx->cb.dst_fn)),
+			 (unsigned int) tree_to_uhwi (purpose));
+		print_generic_expr (dump_file, ovar, dump_flags);
+		fprintf (dump_file, "': kind = %u, size = ",
+			 (unsigned int) tkind);
+		print_generic_expr (dump_file, s, dump_flags);
+		fprintf (dump_file, "\n");
+	      }
+
 	    break;
 	  }
 
       gcc_assert (map_idx == map_cnt);
+      if (flag_checking)
+	for (unsigned int i = 0; i < map_cnt; ++i)
+	  {
+	    tree t_index = (*vsize)[i].index;
+	    unsigned HOST_WIDE_INT ii = tree_to_uhwi (t_index);
+	    gcc_assert (ii == i);
+
+	    t_index = (*vsize)[i].index;
+	    ii = tree_to_uhwi (t_index);
+	    gcc_assert (ii == i);
+	  }
 
       DECL_INITIAL (TREE_VEC_ELT (t, 1))
 	= build_constructor (TREE_TYPE (TREE_VEC_ELT (t, 1)), vsize);
diff --git a/gcc/testsuite/c-c++-common/gomp/lower_omp_target-mappings-1.c b/gcc/testsuite/c-c++-common/gomp/lower_omp_target-mappings-1.c
new file mode 100644
index 000000000000..a3b712e75572
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/lower_omp_target-mappings-1.c
@@ -0,0 +1,50 @@
+// Verify 'lower_omp_target' mappings.
+
+// { dg-additional-options "-fdump-tree-omplower-details" }
+
+int main(int argc, char *argv[])
+{
+  char vla[argc + 3];
+#pragma omp target map(from:vla[0:argc])
+  // { dg-final { scan-tree-dump-times {(?n)^Mapping \[.*lower_omp_target-mappings-1.c:8:9\] main\._omp_fn\.0 \[0\] '\*vla.1\[0\]': kind = 2/15, size = .*, align = 1$} 1 omplower } }
+  // { dg-final { scan-tree-dump-times {(?n)^Mapping \[.*lower_omp_target-mappings-1.c:8:9\] main\._omp_fn\.0 \[1\] '.*': kind = 13, size = 0, align = 4$} 1 omplower { target { ! { c++ && lp64 } } } } }
+  // { dg-final { scan-tree-dump-times {(?n)^Mapping \[.*lower_omp_target-mappings-1.c:8:9\] main\._omp_fn\.0 \[1\] '.*': kind = 13, size = 0, align = 8$} 1 omplower { target { c++ && lp64 } } } }
+  // { dg-final { scan-tree-dump-times {(?n)^Mapping \[.*lower_omp_target-mappings-1.c:8:9\] main\._omp_fn\.0 \[2\] 'argc': kind = 13, size = 0, align = 4$} 1 omplower } }
+  // { dg-final { scan-tree-dump-not {(?n)^Mapping \[.*lower_omp_target-mappings-1.c:8:9\] main\._omp_fn\.0 \[3\]} omplower } }
+  {
+    for (int i = 0; i < argc; ++i)
+      vla[i] = i + sizeof vla;
+  }
+
+#pragma omp target data use_device_ptr(argv)
+  // { dg-final { scan-tree-dump-times {(?n)^Mapping \[.*lower_omp_target-mappings-1.c:19:9\] main \[0\] 'argv': kind = 14, size = 0$} 1 omplower } }
+  // { dg-final { scan-tree-dump-not {(?n)^Mapping \[.*lower_omp_target-mappings-1.c:19:9\] main \[1\]} omplower } }
+  {
+#pragma omp target map(from:argc)
+    // { dg-final { scan-tree-dump-times {(?n)^Mapping \[.*lower_omp_target-mappings-1.c:23:9\] main\._omp_fn\.1 \[0\] 'argc': kind = 2/2, size = 4, align = 4$} 1 omplower } }
+    // { dg-final { scan-tree-dump-times {(?n)^Mapping \[.*lower_omp_target-mappings-1.c:23:9\] main\._omp_fn\.1 \[1\] '\*vla\.1': kind = 3/3, size = .*, align = 1$} 1 omplower } }
+    // { dg-final { scan-tree-dump-times {(?n)^Mapping \[.*lower_omp_target-mappings-1.c:23:9\] main\._omp_fn\.1 \[2\] 'MEM\[\(char \*\)argv\]': kind = 15/15, size = 0, align = 1$} 1 omplower } }
+    // { dg-final { scan-tree-dump-not {(?n)^Mapping \[.*lower_omp_target-mappings-1.c:23:9\] main\._omp_fn\.1 \[3\]} omplower } }
+    {
+      argc = (argv != (void *) vla);
+    }
+  }
+
+#pragma omp target data map(from:argc) 
+  // { dg-final { scan-tree-dump-times {(?n)^Mapping \[.*lower_omp_target-mappings-1.c:33:9\] main \[0\] 'argc': kind = 2/2, size = 4, align = 4$} 1 omplower } }
+  // { dg-final { scan-tree-dump-not {(?n)^Mapping \[.*lower_omp_target-mappings-1.c:33:9\] main \[1\]} omplower } }
+  {
+#pragma omp target is_device_ptr(argv)
+    // { dg-final { scan-tree-dump-times {(?n)^Mapping \[.*lower_omp_target-mappings-1.c:37:9\] main\._omp_fn\.2 \[0\] 'argv': kind = 13, size = 0$} 1 omplower } }
+    // { dg-final { scan-tree-dump-times {(?n)^Mapping \[.*lower_omp_target-mappings-1.c:37:9\] main\._omp_fn\.2 \[1\] '\*vla\.1': kind = 3/3, size = .*, align = 1$} 1 omplower } }
+    // { dg-final { scan-tree-dump-times {(?n)^Mapping \[.*lower_omp_target-mappings-1.c:37:9\] main\._omp_fn\.2 \[2\] 'argc': kind = 13, size = 0, align = 4$} 1 omplower } }
+    // { dg-final { scan-tree-dump-not {(?n)^Mapping \[.*lower_omp_target-mappings-1.c:37:9\] main\._omp_fn\.2 \[3\]} omplower } }
+    {
+      argc = (argv != (void *) &vla[1]);
+    }
+  }
+
+  return 0;
+}
+
+// { dg-final { scan-tree-dump-times {(?n)^Mapping} 11 omplower } }
-- 
2.17.1
Jakub Jelinek May 27, 2019, 7:29 p.m. | #3
On Mon, May 27, 2019 at 09:05:06PM +0200, Thomas Schwinge wrote:
> > The arrays are emitted in the *.omplower dump, so I think it is much better

> > to scan-tree-dump their content

> 

> That's not feasible in the general case.


Why?  The arrays have easily parseable names (.omp_data_kinds.N), then they
have an initializer and you can easily write a regexp to say match 3
occurences of the firstprivate int kind constant in there.

> > if for whatever reason adding a runtime

> > testcase isn't sufficient

> 

> Why should I have an execution test for something that should really be

> verified at the compiler side, at compile time?  And, for example, an


Generally a runtime test can check all compilation phases, linking, runtime,
while what looks like a unit test only tests a small portion of that.
> 

> > over adding further printouts and matching that.

> 

> I had assumed that you'd noticed that a lot of compiler passes are

> dumping stuff that can then be scanned for.  Are you insisting that GCC's


Usually the detailed messages in the dump print why certain changes are
done, or other information that is not readily available in the dump
already.  I'm not convinced what you are doing brings anything that isn't
there already.

	Jakub
Thomas Schwinge May 27, 2019, 9:33 p.m. | #4
Hi!

I'm highly annoyed that we're wasting a lot of time arguing about such a
minor item.

On Mon, 27 May 2019 21:29:13 +0200, Jakub Jelinek <jakub@redhat.com> wrote:
> On Mon, May 27, 2019 at 09:05:06PM +0200, Thomas Schwinge wrote:

> > > The arrays are emitted in the *.omplower dump, so I think it is much better

> > > to scan-tree-dump their content

> > 

> > That's not feasible in the general case.

> 

> Why?  The arrays have easily parseable names (.omp_data_kinds.N), then they

> have an initializer and you can easily write a regexp to say match 3

> occurences of the firstprivate int kind constant in there.


Given the test case that's part of my patch, we get:

    $ grep -F omp_data_kinds < lower_omp_target-mappings-1.c.007t.omplower
              .omp_data_kinds.15[0] = D.2509;
              .omp_data_kinds.15[1] = 781;
              .omp_data_kinds.15[2] = 525;
              #pragma omp target num_teams(1) thread_limit(0) map(from:*vla.1[0] [len: _18]) map(firstprivate:vla [pointer assign, bias: 0]) firstprivate(D.2437) firstprivate(argc) [child fn: main._omp_fn.0 (.omp_data_arr.13, .omp_data_sizes.14, .omp_data_kinds.15)]
              .omp_data_kinds.15 = {CLOBBER};
                            #pragma omp target num_teams(1) thread_limit(0) map(from:argc [len: 4]) map(tofrom:*vla.1 [len: D.2440]) map(firstprivate:vla [pointer assign, bias: 0]) map(alloc:MEM[(char *)argv] [len: 0]) map(firstprivate:argv [pointer assign, bias: 0]) [child fn: main._omp_fn.1 (.omp_data_arr.18, .omp_data_sizes.19, .omp_data_kinds.20)]
                            #pragma omp target num_teams(1) thread_limit(0) is_device_ptr(argv) map(tofrom:*vla.1 [len: D.2440]) map(firstprivate:vla [pointer assign, bias: 0]) firstprivate(argc) [child fn: main._omp_fn.2 (.omp_data_arr.25, .omp_data_sizes.26, .omp_data_kinds.27)]

So that doesn't even display all the arrays.  (The (default) 'static'
ones appear only in the next, 'lower', dump.)

What OMP construct does '.omp_data_kinds.15' relate to?  I can't tell
(what number each one gets), so can have just one per file?

What is 'D.2509'?  I can't follow/scan for that.

The values '781', '525' contain the alignment encoded, which may vary per
GCC target configuration, making it clumsy to scan for.

Etc.

(Why do I even have to explain this to you -- you must aleady before have
come across such dumps, and their peculiarities?)

With my patch, we can instead have:

    // { dg-final { scan-tree-dump-times {(?n)^Mapping \[.*lower_omp_target-mappings-1.c:8:9\] main\._omp_fn\.0 \[0\] '\*vla.1\[0\]': kind = 2/15, size = .*, align = 1$} 1 omplower } }
    // { dg-final { scan-tree-dump-times {(?n)^Mapping \[.*lower_omp_target-mappings-1.c:8:9\] main\._omp_fn\.0 \[1\] '.*': kind = 13, size = 0, align = 4$} 1 omplower { target { ! { c++ && lp64 } } } } }
    // { dg-final { scan-tree-dump-times {(?n)^Mapping \[.*lower_omp_target-mappings-1.c:8:9\] main\._omp_fn\.0 \[1\] '.*': kind = 13, size = 0, align = 8$} 1 omplower { target { c++ && lp64 } } } }
    // { dg-final { scan-tree-dump-times {(?n)^Mapping \[.*lower_omp_target-mappings-1.c:8:9\] main\._omp_fn\.0 \[2\] 'argc': kind = 13, size = 0, align = 4$} 1 omplower } }

> > > if for whatever reason adding a runtime

> > > testcase isn't sufficient

> > 

> > Why should I have an execution test for something that should really be

> > verified at the compiler side, at compile time?  And, for example, an

> 

> Generally a runtime test can check all compilation phases, linking, runtime,

> while what looks like a unit test only tests a small portion of that.


I don't understand what you're intending to explain here.

> > > over adding further printouts and matching that.

> > 

> > I had assumed that you'd noticed that a lot of compiler passes are

> > dumping stuff that can then be scanned for.  Are you insisting that GCC's

> 

> Usually the detailed messages in the dump print why certain changes are

> done, or other information that is not readily available in the dump

> already.


So we must not do something new/useful, because it's... new?

> I'm not convinced what you are doing brings anything that isn't

> there already.


Aha.

Why exactly are you objecting that the compiler produces useful
intermediate dumps, understandable my machine and human?


Grüße
 Thomas
Jakub Jelinek May 28, 2019, 9:21 a.m. | #5
On Mon, May 27, 2019 at 11:33:00PM +0200, Thomas Schwinge wrote:
> Why exactly are you objecting that the compiler produces useful

> intermediate dumps, understandable my machine and human?


While still not convinced we need this, if we add it, you are still
printing a lot of stuff that is useless to humans and makes it harder to
read, plus you unnecessarily print it in two spots rather than just one.

Printing Mapping on each line just adds clutter, just print the target
stmt with TDF_LINENO followed by "Mappings:\n"
once with the location string, replace the break; at the end of
OMP_CLAUSE_FIRSTPRIVATE: handling with nc = c; tkind_zero = tkind; goto
label; which will be right before if (nc && nc != c) and do the printing
there.  print_generic_expr the clause (c), the name of function seems
completely useless to me (and trying to match it in scan-tree-dump when it
can change any time doesn't make sense), TREE_PURPOSE doesn't seem to be
much useful either, print s expression (size), if tkind == tkind_zero,
just print one number (tkind & ((HOST_WIDE_INT_1U << talign_shift) - 1),
otherwise print two numbers extracted similar way from also tkind_zero,
finally the alignment.

	Jakub

Patch

From ba19ef0d5aebf2604e6625d1cca81fc477801966 Mon Sep 17 00:00:00 2001
From: Thomas Schwinge <thomas@codesourcery.com>
Date: Sun, 26 May 2019 18:33:26 +0200
Subject: [PATCH] Make possible 'scan-tree-dump' of 'lower_omp_target' mapping
 kinds

	gcc/
	* omp-low.c (lower_omp_target): Dump mapping kinds.
	gcc/testsuite/
	* c-c++-common/gomp/lower_omp_target-mappings-1.c: New file.
---
 gcc/omp-low.c                                 | 103 +++++++++++++++---
 .../gomp/lower_omp_target-mappings-1.c        |  49 +++++++++
 2 files changed, 136 insertions(+), 16 deletions(-)
 create mode 100644 gcc/testsuite/c-c++-common/gomp/lower_omp_target-mappings-1.c

diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index faab5d384280..79468289d88f 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -9218,15 +9218,23 @@  lower_omp_taskreg (gimple_stmt_iterator *gsi_p, omp_context *ctx)
 static void
 lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
 {
+  pretty_printer pp;
   tree clauses;
   tree child_fn, t, c;
   gomp_target *stmt = as_a <gomp_target *> (gsi_stmt (*gsi_p));
   gbind *tgt_bind, *bind, *dep_bind = NULL;
   gimple_seq tgt_body, olist, ilist, fplist, new_body;
   location_t loc = gimple_location (stmt);
+  const char *loc_str = NULL;
   bool offloaded, data_region;
   unsigned int map_cnt = 0;
 
+  if (dump_file && (dump_flags & TDF_DETAILS))
+    {
+      dump_location (&pp, loc);
+      loc_str = pp_formatted_text (&pp);
+    }
+
   offloaded = is_gimple_omp_offloaded (stmt);
   switch (gimple_omp_target_kind (stmt))
     {
@@ -9687,7 +9695,8 @@  lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
 	    if (TREE_CODE (s) != INTEGER_CST)
 	      TREE_STATIC (TREE_VEC_ELT (t, 1)) = 0;
 
-	    unsigned HOST_WIDE_INT tkind, tkind_zero;
+	    unsigned HOST_WIDE_INT tkind, tkind_align;
+	    unsigned HOST_WIDE_INT tkind_zero, tkind_zero_align;
 	    switch (OMP_CLAUSE_CODE (c))
 	      {
 	      case OMP_CLAUSE_MAP:
@@ -9743,25 +9752,43 @@  lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
 				 < (HOST_WIDE_INT_C (1U) << talign_shift));
 	    gcc_checking_assert (tkind_zero
 				 < (HOST_WIDE_INT_C (1U) << talign_shift));
-	    talign = ceil_log2 (talign);
-	    tkind |= talign << talign_shift;
-	    tkind_zero |= talign << talign_shift;
-	    gcc_checking_assert (tkind
+	    {
+	      unsigned int talign2 = ceil_log2 (talign);
+	      tkind_align = tkind | (talign2 << talign_shift);
+	      tkind_zero_align = tkind_zero | (talign2 << talign_shift);
+	    }
+	    gcc_checking_assert (tkind_align
 				 <= tree_to_uhwi (TYPE_MAX_VALUE (tkind_type)));
-	    gcc_checking_assert (tkind_zero
+	    gcc_checking_assert (tkind_zero_align
 				 <= tree_to_uhwi (TYPE_MAX_VALUE (tkind_type)));
-	    if (tkind == tkind_zero)
-	      x = build_int_cstu (tkind_type, tkind);
+	    if (tkind_align == tkind_zero_align)
+	      x = build_int_cstu (tkind_type, tkind_align);
 	    else
 	      {
 		TREE_STATIC (TREE_VEC_ELT (t, 2)) = 0;
 		x = build3 (COND_EXPR, tkind_type,
 			    fold_build2 (EQ_EXPR, boolean_type_node,
 					 unshare_expr (s), size_zero_node),
-			    build_int_cstu (tkind_type, tkind_zero),
-			    build_int_cstu (tkind_type, tkind));
+			    build_int_cstu (tkind_type, tkind_zero_align),
+			    build_int_cstu (tkind_type, tkind_align));
 	      }
 	    CONSTRUCTOR_APPEND_ELT (vkind, purpose, x);
+
+	    if (dump_file && (dump_flags & TDF_DETAILS))
+	      {
+		fprintf (dump_file, "Mapping %s%s [%u] '",
+			 loc_str,
+			 IDENTIFIER_POINTER (DECL_NAME (ctx->cb.dst_fn)),
+			 (unsigned int) tree_to_uhwi (purpose));
+		print_generic_expr (dump_file, ovar, dump_flags);
+		fprintf (dump_file, "': size = ");
+		print_generic_expr (dump_file, s, dump_flags);
+		fprintf (dump_file, ", kind = %u/%u, align = %u\n",
+			 (unsigned int) tkind,
+			 (unsigned int) tkind_zero,
+			 talign);
+	      }
+
 	    if (nc && nc != c)
 	      c = nc;
 	    break;
@@ -9826,12 +9853,29 @@  lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
 
 	    gcc_checking_assert (tkind
 				 < (HOST_WIDE_INT_C (1U) << talign_shift));
-	    talign = ceil_log2 (talign);
-	    tkind |= talign << talign_shift;
-	    gcc_checking_assert (tkind
+	    {
+	      unsigned int talign2 = ceil_log2 (talign);
+	      tkind_align = tkind | (talign2 << talign_shift);
+	    }
+	    gcc_checking_assert (tkind_align
 				 <= tree_to_uhwi (TYPE_MAX_VALUE (tkind_type)));
 	    CONSTRUCTOR_APPEND_ELT (vkind, purpose,
-				    build_int_cstu (tkind_type, tkind));
+				    build_int_cstu (tkind_type, tkind_align));
+
+	    if (dump_file && (dump_flags & TDF_DETAILS))
+	      {
+		fprintf (dump_file, "Mapping %s%s [%u] '",
+			 loc_str,
+			 IDENTIFIER_POINTER (DECL_NAME (ctx->cb.dst_fn)),
+			 (unsigned int) tree_to_uhwi (purpose));
+		print_generic_expr (dump_file, ovar, dump_flags);
+		fprintf (dump_file, "': size = ");
+		print_generic_expr (dump_file, s, dump_flags);
+		fprintf (dump_file, ", kind = %u, align = %u\n",
+			 (unsigned int) tkind,
+			 talign);
+	      }
+
 	    break;
 
 	  case OMP_CLAUSE_USE_DEVICE_PTR:
@@ -9862,14 +9906,41 @@  lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
 	    CONSTRUCTOR_APPEND_ELT (vsize, purpose, s);
 	    gcc_checking_assert (tkind
 				 < (HOST_WIDE_INT_C (1U) << talign_shift));
-	    gcc_checking_assert (tkind
+	    tkind_align = tkind;
+	    gcc_checking_assert (tkind_align
 				 <= tree_to_uhwi (TYPE_MAX_VALUE (tkind_type)));
 	    CONSTRUCTOR_APPEND_ELT (vkind, purpose,
-				    build_int_cstu (tkind_type, tkind));
+				    build_int_cstu (tkind_type, tkind_align));
+
+	    if (dump_file && (dump_flags & TDF_DETAILS))
+	      {
+		fprintf (dump_file, "Mapping %s%s [%u] '",
+			 loc_str,
+			 IDENTIFIER_POINTER (DECL_NAME (ctx->cb.dst_fn)),
+			 (unsigned int) tree_to_uhwi (purpose));
+		print_generic_expr (dump_file, ovar, dump_flags);
+		fprintf (dump_file, "': size = ");
+		print_generic_expr (dump_file, s, dump_flags);
+		fprintf (dump_file, ", kind = %u, align = %u\n",
+			 (unsigned int) tkind,
+			 talign);
+	      }
+
 	    break;
 	  }
 
       gcc_assert (map_idx == map_cnt);
+      if (flag_checking)
+	for (unsigned int i = 0; i < map_cnt; ++i)
+	  {
+	    tree t_index = (*vsize)[i].index;
+	    unsigned HOST_WIDE_INT ii = tree_to_uhwi (t_index);
+	    gcc_assert (ii == i);
+
+	    t_index = (*vsize)[i].index;
+	    ii = tree_to_uhwi (t_index);
+	    gcc_assert (ii == i);
+	  }
 
       DECL_INITIAL (TREE_VEC_ELT (t, 1))
 	= build_constructor (TREE_TYPE (TREE_VEC_ELT (t, 1)), vsize);
diff --git a/gcc/testsuite/c-c++-common/gomp/lower_omp_target-mappings-1.c b/gcc/testsuite/c-c++-common/gomp/lower_omp_target-mappings-1.c
new file mode 100644
index 000000000000..2e8a3064078f
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/lower_omp_target-mappings-1.c
@@ -0,0 +1,49 @@ 
+// { dg-additional-options "-fdump-tree-omplower-details" }
+
+int main(int argc, char *argv[])
+{
+  char vla[argc + 3];
+#pragma omp target map(from:vla[0:argc])
+  // { dg-final { scan-tree-dump-times {(?n)^Mapping \[.*lower_omp_target-mappings-1.c:6:9\] main\._omp_fn\.0 \[0\] '\*vla.1\[0\]': size = .*, kind = 2/15, align = 1$} 1 omplower } }
+  // { dg-final { scan-tree-dump-times {(?n)^Mapping \[.*lower_omp_target-mappings-1.c:6:9\] main\._omp_fn\.0 \[1\] '.*': size = 0, kind = 13, align = 4$} 1 omplower { target { ! { c++ && lp64 } } } } }
+  // { dg-final { scan-tree-dump-times {(?n)^Mapping \[.*lower_omp_target-mappings-1.c:6:9\] main\._omp_fn\.0 \[1\] '.*': size = 0, kind = 13, align = 8$} 1 omplower { target { c++ && lp64 } } } }
+  // { dg-final { scan-tree-dump-times {(?n)^Mapping \[.*lower_omp_target-mappings-1.c:6:9\] main\._omp_fn\.0 \[2\] 'argc': size = 0, kind = 13, align = 4$} 1 omplower } }
+  // { dg-final { scan-tree-dump-not {(?n)^Mapping \[.*lower_omp_target-mappings-1.c:6:9\] main\._omp_fn\.0 \[3\]} omplower } }
+  {
+    for (int i = 0; i < argc; ++i)
+      vla[i] = i + sizeof vla;
+  }
+
+#pragma omp target data use_device_ptr(argv)
+  // { dg-final { scan-tree-dump-times {(?n)^Mapping \[.*lower_omp_target-mappings-1.c:17:9\] main \[0\] 'argv': size = 0, kind = 14, align = 4$} 1 omplower } }
+  // { dg-final { scan-tree-dump-not {(?n)^Mapping \[.*lower_omp_target-mappings-1.c:17:9\] main \[1\]} omplower } }
+  {
+#pragma omp target map(from:argc)
+    // { dg-final { scan-tree-dump-times {(?n)^Mapping \[.*lower_omp_target-mappings-1.c:21:9\] main\._omp_fn\.1 \[0\] 'argc': size = 4, kind = 2/2, align = 4$} 1 omplower } }
+    // { dg-final { scan-tree-dump-times {(?n)^Mapping \[.*lower_omp_target-mappings-1.c:21:9\] main\._omp_fn\.1 \[1\] '\*vla\.1': size = .*, kind = 3/3, align = 1$} 1 omplower } }
+    // { dg-final { scan-tree-dump-times {(?n)^Mapping \[.*lower_omp_target-mappings-1.c:21:9\] main\._omp_fn\.1 \[2\] 'MEM\[\(char \*\)argv\]': size = 0, kind = 15/15, align = 1$} 1 omplower } }
+    // { dg-final { scan-tree-dump-not {(?n)^Mapping \[.*lower_omp_target-mappings-1.c:21:9\] main\._omp_fn\.1 \[3\]} omplower } }
+    {
+      argc = (argv != (void *) vla);
+    }
+  }
+
+#pragma omp target data map(from:argc) 
+  // { dg-final { scan-tree-dump-times {(?n)^Mapping \[.*lower_omp_target-mappings-1.c:31:9\] main \[0\] 'argc': size = 4, kind = 2/2, align = 4$} 1 omplower } }
+  // { dg-final { scan-tree-dump-not {(?n)^Mapping \[.*lower_omp_target-mappings-1.c:31:9\] main \[1\]} omplower } }
+  {
+#pragma omp target is_device_ptr(argv)
+    // { dg-final { scan-tree-dump-times {(?n)^Mapping \[.*lower_omp_target-mappings-1.c:35:9\] main\._omp_fn\.2 \[0\] 'argv': size = 0, kind = 13, align = 0$} 1 omplower { target c } } }
+    // { dg-final { scan-tree-dump-times {(?n)^Mapping \[.*lower_omp_target-mappings-1.c:35:9\] main\._omp_fn\.2 \[0\] 'argv': size = 0, kind = 13, align = 1$} 1 omplower { target c++ } } }
+    // { dg-final { scan-tree-dump-times {(?n)^Mapping \[.*lower_omp_target-mappings-1.c:35:9\] main\._omp_fn\.2 \[1\] '\*vla\.1': size = .*, kind = 3/3, align = 1$} 1 omplower } }
+    // { dg-final { scan-tree-dump-times {(?n)^Mapping \[.*lower_omp_target-mappings-1.c:35:9\] main\._omp_fn\.2 \[2\] 'argc': size = 0, kind = 13, align = 4$} 1 omplower } }
+    // { dg-final { scan-tree-dump-not {(?n)^Mapping \[.*lower_omp_target-mappings-1.c:35:9\] main\._omp_fn\.2 \[3\]} omplower } }
+    {
+      argc = (argv != (void *) &vla[1]);
+    }
+  }
+
+  return 0;
+}
+
+// { dg-final { scan-tree-dump-times {(?n)^Mapping} 11 omplower } }
-- 
2.17.1