public inbox for gcc-patches@gcc.gnu.org
 help / color / mirror / Atom feed
From: Chung-Lin Tang <cltang@codesourcery.com>
To: gcc-patches <gcc-patches@gcc.gnu.org>,
	Jakub Jelinek <jakub@redhat.com>,
	Tobias Burnus <tobias@codesourcery.com>
Subject: [PATCH, OpenMP] PR103642 - Fix omp-low ICE for indirect references based off component access
Date: Mon, 3 Jan 2022 22:15:26 +0800	[thread overview]
Message-ID: <7233befe-4c44-b275-c4ac-82af937f32f0@codesourcery.com> (raw)

[-- Attachment #1: Type: text/plain, Size: 1100 bytes --]

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.




	

[-- Attachment #2: pr103642.patch --]
[-- Type: text/plain, Size: 1242 bytes --]

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;
+}

             reply	other threads:[~2022-01-03 14:15 UTC|newest]

Thread overview: 3+ messages / expand[flat|nested]  mbox.gz  Atom feed  top
2022-01-03 14:15 Chung-Lin Tang [this message]
2022-01-17 14:55 ` Chung-Lin Tang
2022-01-18 14:09 ` Jakub Jelinek

Reply instructions:

You may reply publicly to this message via plain-text email
using any one of the following methods:

* Save the following mbox file, import it into your mail client,
  and reply-to-all from there: mbox

  Avoid top-posting and favor interleaved quoting:
  https://en.wikipedia.org/wiki/Posting_style#Interleaved_style

* Reply using the --to, --cc, and --in-reply-to
  switches of git-send-email(1):

  git send-email \
    --in-reply-to=7233befe-4c44-b275-c4ac-82af937f32f0@codesourcery.com \
    --to=cltang@codesourcery.com \
    --cc=gcc-patches@gcc.gnu.org \
    --cc=jakub@redhat.com \
    --cc=tobias@codesourcery.com \
    /path/to/YOUR_REPLY

  https://kernel.org/pub/software/scm/git/docs/git-send-email.html

* If your mail client supports setting the In-Reply-To header
  via mailto: links, try the mailto: link
Be sure your reply has a Subject: header at the top and a blank line before the message body.
This is a public inbox, see mirroring instructions
for how to clone and mirror all data and code used for this inbox;
as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).