| Message ID | 8736l1qi3o.fsf@euler.schwinge.homeip.net |
|---|---|
| State | New |
| Headers | show |
| Series | Make possible 'scan-tree-dump' of 'lower_omp_target' mapping kinds | expand |
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
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
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
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
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
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
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