Message ID | de632e08-1867-2388-76e5-565c135ea536@codesourcery.com |
---|---|
State | New |
Headers | show |
Series | [OpenMP] Fix mapping of artificial variables (PR94874) | expand |
On Fri, May 08, 2020 at 05:09:07PM +0200, Tobias Burnus wrote: > I was thinking of simply marking them as > "nflags |= GOVD_FIRSTPRIVATE" but I am not sure whether > that would always make sense, either. In any case, a > simple usage would bypass the > "implicit mapping of assumed size array" > diagnostic in gfc_omp_finish_clause. > > One could also use a value returned by the hook, > but currently it is tailored for shared memory > use only. A fix would be either a new argument > ("bool for_mapping") plus special handling or > a new hook. In any case, the current hook has: I think we want a new hook, the clear cases (mostly DECL_ARTIFICIAL ones, if it is really something compiler created and not something under user's control) handle the way we find best for them (perhaps firstprivate in some cases, though e.g. for typeinfo and vtables I think we want to make them declare target to), but for user-visible stuff best file a github OpenMP/spec issue (can you please do that) and discuss there? I mean, if something is predetermined firstprivate, then one could in some reading assume it then can't be specified and must be firstprivatized on target (or as an exception may be specified but still default to being firstprivatized), but if the spec says that e.g. static data members are predetermined shared, then I don't see how that makes any difference to how it should be mapped or not. So, it is unclear what void bar (const char *); #pragma omp declare target to (bar) void foo () { #pragma omp target defaultmap(none) bar (__FUNCTION__); } should be doing, etc.; target really doesn't have shared clauses and nothing says that variables that have predetermined data sharing of shared are say mapped(tofrom:) implicitly. Jakub
On 5/12/20 1:02 PM, Jakub Jelinek wrote: > I think we want a new hook, the clear cases (mostly DECL_ARTIFICIAL ones, > if it is really something compiler created and not something under user's > control) … Attached is one version, which is somewhat minimalist; I did not check what happens with __FUNCTION__ and also the Fortran part is not very extended but it is a start. Additionally, I fixed the comments for the existing c_omp_predetermined_sharing. OK for the trunk? Tobias PS: Not tested with an actually offloading compiler due to PR ipa/95320. ----------------- Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander Walter
On Mon, May 25, 2020 at 07:10:31PM +0200, Tobias Burnus wrote: > +enum omp_clause_defaultmap_kind > +cxx_omp_predetermined_mapping (tree decl) > +{ > + enum omp_clause_default_kind ret = cxx_omp_predetermined_sharing_1 (decl); > + if (ret != OMP_CLAUSE_DEFAULT_UNSPECIFIED) > + return OMP_CLAUSE_DEFAULTMAP_TO; I don't like the above 3 lines, do you need it for anything? Static data members certainly aren't covered and it is unlikely we want to predetermine them mapped. > + /* Predetermine artificial variables holding integral values, those > + are usually result of gimplify_one_sizepos or SAVE_EXPR > + gimplification. */ > + if (VAR_P (decl) > + && DECL_ARTIFICIAL (decl) > + && INTEGRAL_TYPE_P (TREE_TYPE (decl)) > + && !(DECL_LANG_SPECIFIC (decl) > + && DECL_OMP_PRIVATIZED_MEMBER (decl))) > + return OMP_CLAUSE_DEFAULTMAP_FIRSTPRIVATE; > + > + /* Similarly for typeinfo symbols. */ > + if (VAR_P (decl) && DECL_ARTIFICIAL (decl) && DECL_TINFO_P (decl)) > + return OMP_CLAUSE_DEFAULTMAP_TO; I think the typeinfo symbols and vtables we should implicitly mark declare target to, rather than trying to map them, that will not work properly anyway (they contain embedded function pointers). So I think the above 3 lines should go too. Otherwise LGTM. Jakub
On Tue, May 26, 2020 at 04:06:04PM +0200, Jakub Jelinek wrote: > On Mon, May 25, 2020 at 07:10:31PM +0200, Tobias Burnus wrote: > > +enum omp_clause_defaultmap_kind > > +cxx_omp_predetermined_mapping (tree decl) > > +{ > > + enum omp_clause_default_kind ret = cxx_omp_predetermined_sharing_1 (decl); > > + if (ret != OMP_CLAUSE_DEFAULT_UNSPECIFIED) > > + return OMP_CLAUSE_DEFAULTMAP_TO; > > I don't like the above 3 lines, do you need it for anything? > Static data members certainly aren't covered and it is unlikely we want to > predetermine them mapped. And forgot, for this the behavior is already specified in 5.0 that any reference to this should result in map(tofrom:this[:1]). It is not implemented yet, but we'll need to implement it some other way than through this hook. Jakub
Hi Jakub, Revised patch – changes are: * cp-gimplify.c's cxx_omp_predetermined_mapping (twice 3 lines removed as proposed) * gimplify.c's omp_notice_variable: changes how returned 'kind' is used. Namely: Pre-commit follow-up version to the 'LGTM patch': As instructed, I have removed the two-times three lines from cp-gimplify.c's cxx_omp_predetermined_mapping. However, I made a thinko regarding 'nflags |= kind;' which just happened to work for a test case – the nflags value differs from the enum value. [See gimplify.c's omp_notice_variable; see also line 9610+ for the full set of mappings]. For two items + unreachable, I think 'if' is simpler than 'switch', hence, I used it. Alternatively, one could split-off the case block of ll. 9610+ and use a function call for both. Thoughts? OK? Tobias On 5/26/20 5:43 PM, Jakub Jelinek via Gcc-patches wrote: > On Tue, May 26, 2020 at 04:06:04PM +0200, Jakub Jelinek wrote: >> On Mon, May 25, 2020 at 07:10:31PM +0200, Tobias Burnus wrote: >>> +enum omp_clause_defaultmap_kind >>> +cxx_omp_predetermined_mapping (tree decl) >>> +{ >>> + enum omp_clause_default_kind ret = cxx_omp_predetermined_sharing_1 (decl); >>> + if (ret != OMP_CLAUSE_DEFAULT_UNSPECIFIED) >>> + return OMP_CLAUSE_DEFAULTMAP_TO; >> I don't like the above 3 lines, do you need it for anything? >> Static data members certainly aren't covered and it is unlikely we want to >> predetermine them mapped. > And forgot, for this the behavior is already specified in 5.0 that any > reference to this should result in map(tofrom:this[:1]). > It is not implemented yet, but we'll need to implement it some other way > than through this hook. > > Jakub > ----------------- Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander Walter
On Wed, Jun 03, 2020 at 11:03:50AM +0200, Tobias Burnus wrote: > @@ -2123,6 +2124,27 @@ c_omp_predetermined_sharing (tree decl) > return OMP_CLAUSE_DEFAULT_UNSPECIFIED; > } > > +/* OMP_CLAUSE_DEFAULTMAP_CATEGORY_UNSPECIFIED unless OpenMP mapping attribute > + of DECL is predetermined. */ > + > +enum omp_clause_defaultmap_kind > +c_omp_predetermined_mapping (tree decl) > +{ > + /* Predetermine artificial variables holding integral values, those > + are usually result of gimplify_one_sizepos or SAVE_EXPR > + gimplification. */ > + if (VAR_P (decl) > + && DECL_ARTIFICIAL (decl) > + && INTEGRAL_TYPE_P (TREE_TYPE (decl))) > + return OMP_CLAUSE_DEFAULTMAP_FIRSTPRIVATE; > + > + if (c_omp_predefined_variable (decl)) > + return OMP_CLAUSE_DEFAULTMAP_TO; > + > + return OMP_CLAUSE_DEFAULTMAP_CATEGORY_UNSPECIFIED; > +} ... > @@ -2321,6 +2322,22 @@ cxx_omp_predetermined_sharing (tree decl) > return OMP_CLAUSE_DEFAULT_UNSPECIFIED; > } > > +enum omp_clause_defaultmap_kind > +cxx_omp_predetermined_mapping (tree decl) > +{ > + /* Predetermine artificial variables holding integral values, those > + are usually result of gimplify_one_sizepos or SAVE_EXPR > + gimplification. */ > + if (VAR_P (decl) > + && DECL_ARTIFICIAL (decl) > + && INTEGRAL_TYPE_P (TREE_TYPE (decl)) > + && !(DECL_LANG_SPECIFIC (decl) > + && DECL_OMP_PRIVATIZED_MEMBER (decl))) > + return OMP_CLAUSE_DEFAULTMAP_FIRSTPRIVATE; I'd prefer consistency between C and C++, so I'd add + if (c_omp_predefined_variable (decl)) + return OMP_CLAUSE_DEFAULTMAP_TO; here too for now (and I'll create a new spec issue for __FUNCTION__ etc. mapping because the old one happened to be closed without actually solving what is needed). > + > + return OMP_CLAUSE_DEFAULTMAP_CATEGORY_UNSPECIFIED; > +} Ok with that change. Thanks. Jakub
[OpenMP] Fix mapping of artificial variables (PR94874) gcc/ 2020-05-08 Tobias Burnus <tobias@codesourcery.com> PR middle-end/94874 * gimplify.c (omp_notice_variable): For mapping, also call lang_hooks.decls.omp_predetermined_sharing. gcc/testsuite/ 2020-05-08 Thomas Schwinge <thomas@codesourcery.com> Tobias Burnus <tobias@codesourcery.com> PR middle-end/94874 * c-c++-common/gomp/pr94874.c: New. gcc/gimplify.c | 5 ++++- gcc/testsuite/c-c++-common/gomp/pr94874.c | 27 +++++++++++++++++++++++++++ 2 files changed, 31 insertions(+), 1 deletion(-) diff --git a/gcc/gimplify.c b/gcc/gimplify.c index 1d532e6f373..4c5eb0b3d54 100644 --- a/gcc/gimplify.c +++ b/gcc/gimplify.c @@ -7441,7 +7441,10 @@ omp_notice_variable (struct gimplify_omp_ctx *ctx, tree decl, bool in_code) gdmk = GDMK_SCALAR; else gdmk = GDMK_AGGREGATE; - if (ctx->defaultmap[gdmk] == 0) + if (lang_hooks.decls.omp_predetermined_sharing (decl) + != OMP_CLAUSE_DEFAULT_UNSPECIFIED) + ; /* Handled by lang_hooks.decls.omp_finish_clause. */ + else if (ctx->defaultmap[gdmk] == 0) { tree d = lang_hooks.decls.omp_report_decl (decl); error ("%qE not specified in enclosing %<target%>", diff --git a/gcc/testsuite/c-c++-common/gomp/pr94874.c b/gcc/testsuite/c-c++-common/gomp/pr94874.c new file mode 100644 index 00000000000..36da2471a80 --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/pr94874.c @@ -0,0 +1,27 @@ +/* { dg-do compile } */ +/* { dg-additional-options "-fdump-tree-gimple" } */ + +#include <stddef.h> + +size_t +vla (int array_li) +{ + float array[array_li]; + size_t size1, size2; + +#pragma omp parallel default(none) shared(size1, array) + size1 = sizeof array; + +#pragma omp target defaultmap(none) map(from:size2) map(alloc:array) + size2 = sizeof array; + + return size1 + size2; +} + +/* C */ +/* { dg-final { scan-tree-dump "omp parallel .*shared\\(array_li\.\[0-9\]\\)" "gimple" { target { ! c++ } } } } */ +/* { dg-final { scan-tree-dump "omp target .*map\\(tofrom:array_li\.\[0-9\] " "gimple" { target { ! c++ } } } } */ + +/* C++ */ +/* { dg-final { scan-tree-dump "omp parallel .*shared\\(D\.\[0-9\]*\\)" "gimple" { target { c++ } } } } */ +/* { dg-final { scan-tree-dump "omp target .*map\\(tofrom:D\.\[0-9\]* " "gimple" { target { c++ } } } } */