Hi!
On Mon, 27 May 2019 18:51:22 +0200, Jakub Jelinek <ja...@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 <tho...@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
