public inbox for gcc-patches@gcc.gnu.org
 help / color / mirror / Atom feed
From: Chung-Lin Tang <chunglin.tang@siemens.com>
To: gcc-patches <gcc-patches@gcc.gnu.org>,
	Thomas Schwinge <thomas@codesourcery.com>,
	Catherine Moore <clm@codesourcery.com>
Subject: [PATCH, OpenACC 2.7] Implement default clause support for data constructs
Date: Tue, 6 Jun 2023 23:11:55 +0800	[thread overview]
Message-ID: <f2b7a229-0ec8-5404-fea2-4612ffb73bf2@siemens.com> (raw)

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

Hi Thomas,
this patch implements the OpenACC 2.7 addition of default(none|present) support
for data constructs.

Apart from adjusting the front-ends for allowed clauses masks (for acc data),
mostly implemented in gimplify.

Tested on powerpc64le-linux/nvptx, x86_64-linux/amdgcn tests in progress (expect
no surprises). Is this okay for trunk?

Thanks,
Chung-Lin
gcc/c/ChangeLog:

	* c-parser.cc (OACC_DATA_CLAUSE_MASK): Add PRAGMA_OACC_CLAUSE_DEFAULT.

gcc/cp/ChangeLog:

	* parser.cc (OACC_DATA_CLAUSE_MASK): Add PRAGMA_OACC_CLAUSE_DEFAULT.

gcc/fortran/ChangeLog:

	* openmp.cc (OACC_DATA_CLAUSES): Add OMP_CLAUSE_DEFAULT.

gcc/ChangeLog:

	* gimplify.cc (struct gimplify_omp_ctx): Add oacc_data_default_kind
	field.
	(new_omp_context): Initialize oacc_data_default_kind field.
	(gimplify_scan_omp_clauses): Set oacc_data_default_kind for data
	constructs. Set ctx->default_kind for compute constructs from
	ctx->oacc_data_default_kind.

gcc/testsuite/ChangeLog:

	* c-c++-common/goacc/default-3.c: Adjust testcase.
	* c-c++-common/goacc/default-5.c: Adjust testcase.
	* gfortran.dg/goacc/default-3.f95: Adjust testcase.
	* gfortran.dg/goacc/default-5.f: Adjust testcase.

[-- Attachment #2: 0002-OpenACC-2.7-default-clause-support-for-data-construc.patch --]
[-- Type: text/plain, Size: 11311 bytes --]

From 101305aee9b27c6df00d7c403e469bdf8d7f45a4 Mon Sep 17 00:00:00 2001
From: Chung-Lin Tang <cltang@codesourcery.com>
Date: Tue, 6 Jun 2023 03:46:29 -0700
Subject: [PATCH 2/2] OpenACC 2.7: default clause support for data constructs

This patch implements the OpenACC 2.7 addition of default(none|present) support
for data constructs.

Now, specifying "default(none|present)" on a data construct turns on same
default clause behavior for all enclosed compute constructs (which don't
already themselves have a default clause).

gcc/c/ChangeLog:

	* c-parser.cc (OACC_DATA_CLAUSE_MASK): Add PRAGMA_OACC_CLAUSE_DEFAULT.

gcc/cp/ChangeLog:

	* parser.cc (OACC_DATA_CLAUSE_MASK): Add PRAGMA_OACC_CLAUSE_DEFAULT.

gcc/fortran/ChangeLog:

	* openmp.cc (OACC_DATA_CLAUSES): Add OMP_CLAUSE_DEFAULT.

gcc/ChangeLog:

	* gimplify.cc (struct gimplify_omp_ctx): Add oacc_data_default_kind
	field.
	(new_omp_context): Initialize oacc_data_default_kind field.
	(gimplify_scan_omp_clauses): Set oacc_data_default_kind for data
	constructs. Set ctx->default_kind for compute constructs from
	ctx->oacc_data_default_kind.

gcc/testsuite/ChangeLog:

	* c-c++-common/goacc/default-3.c: Adjust testcase.
	* c-c++-common/goacc/default-5.c: Adjust testcase.
	* gfortran.dg/goacc/default-3.f95: Adjust testcase.
	* gfortran.dg/goacc/default-5.f: Adjust testcase.
---
 gcc/c/c-parser.cc                             |  1 +
 gcc/cp/parser.cc                              |  1 +
 gcc/fortran/openmp.cc                         |  3 ++-
 gcc/gimplify.cc                               | 20 +++++++++++++++++++
 gcc/testsuite/c-c++-common/goacc/default-3.c  | 15 +++++++++++++-
 gcc/testsuite/c-c++-common/goacc/default-5.c  | 18 +++++++++++++++--
 gcc/testsuite/gfortran.dg/goacc/default-3.f95 | 15 ++++++++++++++
 gcc/testsuite/gfortran.dg/goacc/default-5.f   | 17 ++++++++++++++--
 8 files changed, 84 insertions(+), 6 deletions(-)

diff --git a/gcc/c/c-parser.cc b/gcc/c/c-parser.cc
index b61aef8b1a2..645d28b320d 100644
--- a/gcc/c/c-parser.cc
+++ b/gcc/c/c-parser.cc
@@ -18133,6 +18133,7 @@ c_parser_oacc_cache (location_t loc, c_parser *parser)
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYIN)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYOUT)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_CREATE)		\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEFAULT)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICEPTR)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF)			\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NO_CREATE)		\
diff --git a/gcc/cp/parser.cc b/gcc/cp/parser.cc
index dd7638f1c93..4b4df29a406 100644
--- a/gcc/cp/parser.cc
+++ b/gcc/cp/parser.cc
@@ -45759,6 +45759,7 @@ cp_parser_oacc_cache (cp_parser *parser, cp_token *pragma_tok)
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYIN)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYOUT)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_CREATE)		\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEFAULT)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DETACH)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICEPTR)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF)			\
diff --git a/gcc/fortran/openmp.cc b/gcc/fortran/openmp.cc
index 4c30548567f..b785e71f20f 100644
--- a/gcc/fortran/openmp.cc
+++ b/gcc/fortran/openmp.cc
@@ -3645,7 +3645,8 @@ error:
 #define OACC_DATA_CLAUSES \
   (omp_mask (OMP_CLAUSE_IF) | OMP_CLAUSE_DEVICEPTR  | OMP_CLAUSE_COPY	      \
    | OMP_CLAUSE_COPYIN | OMP_CLAUSE_COPYOUT | OMP_CLAUSE_CREATE		      \
-   | OMP_CLAUSE_NO_CREATE | OMP_CLAUSE_PRESENT | OMP_CLAUSE_ATTACH)
+   | OMP_CLAUSE_NO_CREATE | OMP_CLAUSE_PRESENT | OMP_CLAUSE_ATTACH	      \
+   | OMP_CLAUSE_DEFAULT)
 #define OACC_LOOP_CLAUSES \
   (omp_mask (OMP_CLAUSE_COLLAPSE) | OMP_CLAUSE_GANG | OMP_CLAUSE_WORKER	      \
    | OMP_CLAUSE_VECTOR | OMP_CLAUSE_SEQ | OMP_CLAUSE_INDEPENDENT	      \
diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc
index 4aa6229fc74..368f5fd7ec8 100644
--- a/gcc/gimplify.cc
+++ b/gcc/gimplify.cc
@@ -225,6 +225,7 @@ struct gimplify_omp_ctx
   vec<tree> loop_iter_var;
   location_t location;
   enum omp_clause_default_kind default_kind;
+  enum omp_clause_default_kind oacc_data_default_kind;
   enum omp_region_type region_type;
   enum tree_code code;
   bool combined_loop;
@@ -459,6 +460,8 @@ new_omp_context (enum omp_region_type region_type)
     c->default_kind = OMP_CLAUSE_DEFAULT_SHARED;
   else
     c->default_kind = OMP_CLAUSE_DEFAULT_UNSPECIFIED;
+  if (gimplify_omp_ctxp)
+    c->oacc_data_default_kind = gimplify_omp_ctxp->oacc_data_default_kind;
   c->defaultmap[GDMK_SCALAR] = GOVD_MAP;
   c->defaultmap[GDMK_SCALAR_TARGET] = GOVD_MAP;
   c->defaultmap[GDMK_AGGREGATE] = GOVD_MAP;
@@ -12050,6 +12053,8 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 
 	case OMP_CLAUSE_DEFAULT:
 	  ctx->default_kind = OMP_CLAUSE_DEFAULT_KIND (c);
+	  if (code == OACC_DATA)
+	    ctx->oacc_data_default_kind = OMP_CLAUSE_DEFAULT_KIND (c);
 	  break;
 
 	case OMP_CLAUSE_INCLUSIVE:
@@ -12098,6 +12103,21 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 	list_p = &OMP_CLAUSE_CHAIN (c);
     }
 
+  if ((code == OACC_PARALLEL
+       || code == OACC_KERNELS
+       || code == OACC_SERIAL)
+      && ctx->default_kind == OMP_CLAUSE_DEFAULT_SHARED
+      && ctx->oacc_data_default_kind != OMP_CLAUSE_DEFAULT_UNSPECIFIED)
+    {
+      ctx->default_kind = ctx->oacc_data_default_kind;
+
+      /* Append actual default clause on compute construct. Not really needed
+	 for omp_notice_variable to work properly, but for debug dump files.  */
+      c = build_omp_clause (UNKNOWN_LOCATION, OMP_CLAUSE_DEFAULT);
+      OMP_CLAUSE_DEFAULT_KIND (c) = ctx->oacc_data_default_kind;
+      *list_p = c;
+    }
+
   ctx->clauses = *orig_list_p;
   gimplify_omp_ctxp = ctx;
 }
diff --git a/gcc/testsuite/c-c++-common/goacc/default-3.c b/gcc/testsuite/c-c++-common/goacc/default-3.c
index ac169a903e9..060629057c1 100644
--- a/gcc/testsuite/c-c++-common/goacc/default-3.c
+++ b/gcc/testsuite/c-c++-common/goacc/default-3.c
@@ -4,7 +4,7 @@ void f1 ()
 {
   int f1_a = 2;
   float f1_b[2];
-  
+
 #pragma acc kernels default (none) /* { dg-message "enclosing OpenACC .kernels. construct" } */
   {
     f1_b[0] /* { dg-error ".f1_b. not specified in enclosing OpenACC .kernels. construct" } */
@@ -15,4 +15,17 @@ void f1 ()
     f1_b[0] /* { dg-error ".f1_b. not specified in enclosing OpenACC .parallel. construct" } */
       = f1_a; /* { dg-error ".f1_a. not specified in enclosing OpenACC .parallel. construct" } */
   }
+
+#pragma acc data default (none)
+#pragma acc kernels /* { dg-message "enclosing OpenACC .kernels. construct" } */
+  {
+    f1_b[0] /* { dg-error ".f1_b. not specified in enclosing OpenACC .kernels. construct" } */
+      = f1_a; /* { dg-error ".f1_a. not specified in enclosing OpenACC .kernels. construct" } */
+  }
+#pragma acc data default (none)
+#pragma acc parallel /* { dg-message "enclosing OpenACC .parallel. construct" } */
+  {
+    f1_b[0] /* { dg-error ".f1_b. not specified in enclosing OpenACC .parallel. construct" } */
+      = f1_a; /* { dg-error ".f1_a. not specified in enclosing OpenACC .parallel. construct" } */
+  }
 }
diff --git a/gcc/testsuite/c-c++-common/goacc/default-5.c b/gcc/testsuite/c-c++-common/goacc/default-5.c
index 37e3c3555cd..e44a0dd49d7 100644
--- a/gcc/testsuite/c-c++-common/goacc/default-5.c
+++ b/gcc/testsuite/c-c++-common/goacc/default-5.c
@@ -4,8 +4,8 @@
 
 void f1 ()
 {
-  int f1_a = 2;
-  float f1_b[2];
+  int f1_a = 2, f1_c = 3;
+  float f1_b[2], f1_d[2];
 
 #pragma acc kernels default (present)
   /* { dg-final { scan-tree-dump-times "omp target oacc_kernels default\\(present\\) map\\(force_present:f1_b \[^\\)\]+\\) map\\(force_tofrom:f1_a" 1 "gimple" } } */
@@ -17,4 +17,18 @@ void f1 ()
   {
     f1_b[0] = f1_a;
   }
+
+  /* { dg-final { scan-tree-dump-times "omp target oacc_data default\\(present\\)" 2 "gimple" } } */
+#pragma acc data default (present)
+#pragma acc kernels
+  /* { dg-final { scan-tree-dump-times "omp target oacc_kernels default\\(present\\) map\\(force_present:f1_d \[^\\)\]+\\) map\\(force_tofrom:f1_c" 1 "gimple" } } */
+  {
+    f1_d[0] = f1_c;
+  }
+#pragma acc data default (present)
+#pragma acc parallel
+  /* { dg-final { scan-tree-dump-times "omp target oacc_parallel default\\(present\\) map\\(force_present:f1_d \[^\\)\]+\\) firstprivate\\(f1_c\\)" 1 "gimple" } } */
+  {
+    f1_d[0] = f1_c;
+  }
 }
diff --git a/gcc/testsuite/gfortran.dg/goacc/default-3.f95 b/gcc/testsuite/gfortran.dg/goacc/default-3.f95
index 98ed34200c6..13ccce30fd7 100644
--- a/gcc/testsuite/gfortran.dg/goacc/default-3.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/default-3.f95
@@ -15,4 +15,19 @@ subroutine f1
        = f1_a; ! { dg-error ".f1_a. not specified in enclosing OpenACC .parallel. construct" }
   ! { dg-bogus ".f1_b. not specified in enclosing OpenACC .parallel. construct" "" { xfail *-*-* } .-1 }
   !$acc end parallel
+
+  !$acc data default (none)
+  !$acc kernels ! { dg-message "enclosing OpenACC .kernels. construct" }
+  f1_b(1) & ! { dg-error ".f1_b. not specified in enclosing OpenACC .kernels. construct" "" { xfail *-*-* } }
+       = f1_a; ! { dg-error ".f1_a. not specified in enclosing OpenACC .kernels. construct" }
+  ! { dg-bogus ".f1_b. not specified in enclosing OpenACC .kernels. construct" "" { xfail *-*-* } .-1 }
+  !$acc end kernels
+  !$acc end data
+  !$acc data default (none)
+  !$acc parallel ! { dg-message "enclosing OpenACC .parallel. construct" }
+  f1_b(1) & ! { dg-error ".f1_b. not specified in enclosing OpenACC .parallel. construct" "" { xfail *-*-* } }
+       = f1_a; ! { dg-error ".f1_a. not specified in enclosing OpenACC .parallel. construct" }
+  ! { dg-bogus ".f1_b. not specified in enclosing OpenACC .parallel. construct" "" { xfail *-*-* } .-1 }
+  !$acc end parallel
+  !$acc end data
 end subroutine f1
diff --git a/gcc/testsuite/gfortran.dg/goacc/default-5.f b/gcc/testsuite/gfortran.dg/goacc/default-5.f
index 9dc83cbe601..4424f9e7523 100644
--- a/gcc/testsuite/gfortran.dg/goacc/default-5.f
+++ b/gcc/testsuite/gfortran.dg/goacc/default-5.f
@@ -4,8 +4,8 @@
 
       SUBROUTINE F1
       IMPLICIT NONE
-      INTEGER :: F1_A = 2
-      REAL, DIMENSION (2) :: F1_B
+      INTEGER :: F1_A = 2, F1_C = 3
+      REAL, DIMENSION (2) :: F1_B, F1_D
 
 !$ACC KERNELS DEFAULT (PRESENT)
 ! { dg-final { scan-tree-dump-times "omp target oacc_kernels default\\(present\\) map\\(force_present:f1_b \[^\\)\]+\\) map\\(force_tofrom:f1_a" 1 "gimple" } }
@@ -15,4 +15,17 @@
 ! { dg-final { scan-tree-dump-times "omp target oacc_parallel default\\(present\\) map\\(force_present:f1_b \[^\\)\]+\\) firstprivate\\(f1_a\\)" 1 "gimple" } }
       F1_B(1) = F1_A;
 !$ACC END PARALLEL
+
+!$ACC DATA DEFAULT (PRESENT)
+!$ACC KERNELS
+! { dg-final { scan-tree-dump-times "omp target oacc_kernels default\\(present\\) map\\(force_present:f1_d \[^\\)\]+\\) map\\(force_tofrom:f1_c" 1 "gimple" } }
+      F1_D(1) = F1_C;
+!$ACC END KERNELS
+!$ACC END DATA
+!$ACC DATA DEFAULT (PRESENT)
+!$ACC PARALLEL DEFAULT (PRESENT)
+! { dg-final { scan-tree-dump-times "omp target oacc_parallel default\\(present\\) map\\(force_present:f1_d \[^\\)\]+\\) firstprivate\\(f1_c\\)" 1 "gimple" } }
+      F1_D(1) = F1_C;
+!$ACC END PARALLEL
+!$ACC END DATA
       END SUBROUTINE F1
-- 
2.27.0


             reply	other threads:[~2023-06-06 15:12 UTC|newest]

Thread overview: 5+ messages / expand[flat|nested]  mbox.gz  Atom feed  top
2023-06-06 15:11 Chung-Lin Tang [this message]
2023-06-23 10:47 ` Thomas Schwinge
2023-07-14 10:34   ` Chung-Lin Tang
2023-08-15 14:35     ` [v3] OpenACC 2.7: default clause support for data constructs (was: [PATCH, OpenACC 2.7, v2] Implement default clause support for data constructs) Thomas Schwinge
2023-08-01 15:35 [PATCH, OpenACC 2.7, v2] Implement default clause support for data constructs Chung-Lin Tang

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=f2b7a229-0ec8-5404-fea2-4612ffb73bf2@siemens.com \
    --to=chunglin.tang@siemens.com \
    --cc=clm@codesourcery.com \
    --cc=gcc-patches@gcc.gnu.org \
    --cc=thomas@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).