| Message ID | 7233befe-4c44-b275-c4ac-82af937f32f0@codesourcery.com |
|---|---|
| State | New |
| Headers | show |
| Series | [OpenMP] PR103642 - Fix omp-low ICE for indirect references based off component access | expand |
Ping. On 2022/1/3 10:15 PM, Chung-Lin Tang wrote: > This issue was triggered after the patch extending syntax for component access > in map clauses > (https://gcc.gnu.org/git/gitweb.cgi?p=gcc.git;h=0ab29cf0bb68960c) > > In gimplify_scan_omp_clauses, the case for handling indirect accesses (which creates > firstprivate ptr and zero-length array section map for such decls) was erroneously > went into for non-pointer cases (here being the base struct decl), so added the > appropriate checks there. > > Added new testcase is a compile only test for the ICE. The original omptests t-partial-struct > test actually should not execute correctly, because for map(t.s->a[:N]), map(t.s[:1]) > is not implicitly mapped, thus the entire offloaded access does not work as is. > (fixing that omptests test is out of scope here) > > Tested without regressions, okay for trunk? > > Thanks, > Chung-Lin > > 2022-01-03 Chung-Lin Tang <cltang@codesourcery.com> > > gcc/ChangeLog: > > PR middle-end/103642 > * gimplify.c (gimplify_scan_omp_clauses): Do not do indir_p handling > for non-pointer or non-reference-to-pointer cases. > > gcc/testsuite/ChangeLog: > > * c-c++-common/gomp/pr103642.c: New test. > > > > >
On Mon, Jan 03, 2022 at 10:15:26PM +0800, Chung-Lin Tang wrote: > This issue was triggered after the patch extending syntax for component access > in map clauses > (https://gcc.gnu.org/git/gitweb.cgi?p=gcc.git;h=0ab29cf0bb68960c) > > In gimplify_scan_omp_clauses, the case for handling indirect accesses (which creates > firstprivate ptr and zero-length array section map for such decls) was erroneously > went into for non-pointer cases (here being the base struct decl), so added the > appropriate checks there. > > Added new testcase is a compile only test for the ICE. The original omptests t-partial-struct > test actually should not execute correctly, because for map(t.s->a[:N]), map(t.s[:1]) > is not implicitly mapped, thus the entire offloaded access does not work as is. > (fixing that omptests test is out of scope here) > > Tested without regressions, okay for trunk? > > Thanks, > Chung-Lin > > 2022-01-03 Chung-Lin Tang <cltang@codesourcery.com> > > gcc/ChangeLog: > > PR middle-end/103642 > * gimplify.c (gimplify_scan_omp_clauses): Do not do indir_p handling > for non-pointer or non-reference-to-pointer cases. > > gcc/testsuite/ChangeLog: > > * c-c++-common/gomp/pr103642.c: New test. Ok, with a small nit: > +int main (void) > +{ > + T t; > + t.s = (S *) malloc (sizeof (S)); > + t.s->a = (int *) malloc (sizeof(int) * N); > + > + #pragma omp target map(from: t.s->a[:N]) > + { > + t.s->a[0] = 1; > + } Please free the pointers afterwards. Jakub
diff --git a/gcc/gimplify.c b/gcc/gimplify.c index b118c72f62c..bdc8189c2a7 100644 --- a/gcc/gimplify.c +++ b/gcc/gimplify.c @@ -9543,7 +9543,10 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, == REFERENCE_TYPE)) decl = TREE_OPERAND (decl, 0); } - if (decl != orig_decl && DECL_P (decl) && indir_p) + if (decl != orig_decl && DECL_P (decl) && indir_p + && (TREE_CODE (TREE_TYPE (decl)) == POINTER_TYPE + || (decl_ref + && TREE_CODE (TREE_TYPE (decl_ref)) == POINTER_TYPE))) { gomp_map_kind k = ((code == OACC_EXIT_DATA || code == OMP_TARGET_EXIT_DATA) diff --git a/gcc/testsuite/c-c++-common/gomp/pr103642.c b/gcc/testsuite/c-c++-common/gomp/pr103642.c new file mode 100644 index 00000000000..c5451596b69 --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/pr103642.c @@ -0,0 +1,31 @@ +/* PR middle-end/103642 */ +/* { dg-do compile } */ + +#include <stdlib.h> + +typedef struct +{ + int *a; +} S; + +typedef struct +{ + S *s; + int *ptr; +} T; + +#define N 10 + +int main (void) +{ + T t; + t.s = (S *) malloc (sizeof (S)); + t.s->a = (int *) malloc (sizeof(int) * N); + + #pragma omp target map(from: t.s->a[:N]) + { + t.s->a[0] = 1; + } + + return 0; +}