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.
From 101305aee9b27c6df00d7c403e469bdf8d7f45a4 Mon Sep 17 00:00:00 2001
From: Chung-Lin Tang <clt...@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 via email to