[OpenMP] PR103642 - Fix omp-low ICE for indirect references based off component access
Commit Message
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.
Comments
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
@@ -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)
new file mode 100644
@@ -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;
+}