This is the mail archive of the gcc-patches@gcc.gnu.org mailing list for the GCC project.
| Index Nav: | [Date Index] [Subject Index] [Author Index] [Thread Index] | |
|---|---|---|
| Message Nav: | [Date Prev] [Date Next] | [Thread Prev] [Thread Next] |
| Other format: | [Raw text] | |
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
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
Attachment:
signature.asc
Description: PGP signature
| Index Nav: | [Date Index] [Subject Index] [Author Index] [Thread Index] | |
|---|---|---|
| Message Nav: | [Date Prev] [Date Next] | [Thread Prev] [Thread Next] |