diff mbox series

[OpenACC,2.7,v2] Implement default clause support for data constructs

Message ID 3f274de5-cc52-39c7-c399-f85a1a1a4640@siemens.com
State New
Headers show
Series [OpenACC,2.7,v2] Implement default clause support for data constructs | expand

Commit Message

Chung-Lin Tang Aug. 1, 2023, 3:35 p.m. UTC 
Hi Thomas,
this is v2 of the patch for implementing the OpenACC 2.7 addition of
default(none|present) support for data constructs.

Instead of propagating an additional 'oacc_default_kind' for OpenACC,
this patch does it in a more complete way: it directly propagates the
gimplify_omp_ctx* pointer of the inner most context where we found
a default-clause. This supports displaying the location/type of OpenACC
construct where the default-clause is in the error messages.

The testcases also have the multiple nested data construct testing added,
where we can now have messages referring precisely to the exact innermost
default clause that was active at that program point.

Note, I got rid of the dummy OMP_CLAUSE_DEFAULT creation in this version,
since it seemed not really needed.

Re-tested on master on powerpc64le-linux/nvptx. Okay to commit?

Thanks,
Chung-Lin

2023-08-01  Chung-Lin Tang  <cltang@codesourcery.com>

	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_default_clause_ctx
	field.
	(new_omp_context): Initialize oacc_default_clause_ctx field.
	(oacc_region_type_name): New function.
	(oacc_default_clause): Lookup current default_kind value from
	ctx->oacc_default_clause_ctx, adjust default(none) error and inform
	message dumping.
	(gimplify_scan_omp_clauses): Upon OMP_CLAUSE_DEFAULT case, set
	ctx->oacc_default_clause_ctx to current context.

	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.
diff mbox series

Patch

diff --git a/gcc/c/c-parser.cc b/gcc/c/c-parser.cc
index 24a6eb6e459..974f0132787 100644
--- a/gcc/c/c-parser.cc
+++ b/gcc/c/c-parser.cc
@@ -18196,6 +18196,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 d7ef5b34d42..bc59fbeac20 100644
--- a/gcc/cp/parser.cc
+++ b/gcc/cp/parser.cc
@@ -45860,6 +45860,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 2952cd300ac..c37f843ec3b 100644
--- a/gcc/fortran/openmp.cc
+++ b/gcc/fortran/openmp.cc
@@ -3802,7 +3802,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 320920ed74c..ec0ccc67da8 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;
+  struct gimplify_omp_ctx *oacc_default_clause_ctx;
   enum omp_region_type region_type;
   enum tree_code code;
   bool combined_loop;
@@ -459,6 +460,10 @@  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_default_clause_ctx = gimplify_omp_ctxp->oacc_default_clause_ctx;
+  else
+    c->oacc_default_clause_ctx = c;
   c->defaultmap[GDMK_SCALAR] = GOVD_MAP;
   c->defaultmap[GDMK_SCALAR_TARGET] = GOVD_MAP;
   c->defaultmap[GDMK_AGGREGATE] = GOVD_MAP;
@@ -7699,6 +7704,25 @@  omp_default_clause (struct gimplify_omp_ctx *ctx, tree decl,
   return flags;
 }
 
+/* Return string name for types of OpenACC constructs from ORT_* values.  */
+
+static const char *
+oacc_region_type_name (enum omp_region_type region_type)
+{
+  switch (region_type)
+    {
+    case ORT_ACC_DATA:
+      return "data";
+    case ORT_ACC_PARALLEL:
+      return "parallel";
+    case ORT_ACC_KERNELS:
+      return "kernels";
+    case ORT_ACC_SERIAL:
+      return "serial";
+    default:
+      gcc_unreachable ();
+    }
+}
 
 /* Determine outer default flags for DECL mentioned in an OACC region
    but not declared in an enclosing clause.  */
@@ -7706,7 +7730,6 @@  omp_default_clause (struct gimplify_omp_ctx *ctx, tree decl,
 static unsigned
 oacc_default_clause (struct gimplify_omp_ctx *ctx, tree decl, unsigned flags)
 {
-  const char *rkind;
   bool on_device = false;
   bool is_private = false;
   bool declared = is_oacc_declared (decl);
@@ -7735,17 +7758,20 @@  oacc_default_clause (struct gimplify_omp_ctx *ctx, tree decl, unsigned flags)
       flags |= GOVD_MAP_TO_ONLY;
     }
 
+  /* Use the enclosing construct with a default clause to set the current
+     default kind.  */
+  enum omp_clause_default_kind default_kind
+    = ctx->oacc_default_clause_ctx->default_kind;
+
   switch (ctx->region_type)
     {
     case ORT_ACC_KERNELS:
-      rkind = "kernels";
-
       if (is_private)
 	flags |= GOVD_FIRSTPRIVATE;
       else if (AGGREGATE_TYPE_P (type))
 	{
 	  /* Aggregates default to 'present_or_copy', or 'present'.  */
-	  if (ctx->default_kind != OMP_CLAUSE_DEFAULT_PRESENT)
+	  if (default_kind != OMP_CLAUSE_DEFAULT_PRESENT)
 	    flags |= GOVD_MAP;
 	  else
 	    flags |= GOVD_MAP | GOVD_MAP_FORCE_PRESENT;
@@ -7758,8 +7784,6 @@  oacc_default_clause (struct gimplify_omp_ctx *ctx, tree decl, unsigned flags)
 
     case ORT_ACC_PARALLEL:
     case ORT_ACC_SERIAL:
-      rkind = ctx->region_type == ORT_ACC_PARALLEL ? "parallel" : "serial";
-
       if (is_private)
 	flags |= GOVD_FIRSTPRIVATE;
       else if (on_device || declared)
@@ -7767,7 +7791,7 @@  oacc_default_clause (struct gimplify_omp_ctx *ctx, tree decl, unsigned flags)
       else if (AGGREGATE_TYPE_P (type))
 	{
 	  /* Aggregates default to 'present_or_copy', or 'present'.  */
-	  if (ctx->default_kind != OMP_CLAUSE_DEFAULT_PRESENT)
+	  if (default_kind != OMP_CLAUSE_DEFAULT_PRESENT)
 	    flags |= GOVD_MAP;
 	  else
 	    flags |= GOVD_MAP | GOVD_MAP_FORCE_PRESENT;
@@ -7785,16 +7809,20 @@  oacc_default_clause (struct gimplify_omp_ctx *ctx, tree decl, unsigned flags)
   if (DECL_ARTIFICIAL (decl))
     ; /* We can get compiler-generated decls, and should not complain
 	 about them.  */
-  else if (ctx->default_kind == OMP_CLAUSE_DEFAULT_NONE)
+  else if (default_kind == OMP_CLAUSE_DEFAULT_NONE)
     {
       error ("%qE not specified in enclosing OpenACC %qs construct",
-	     DECL_NAME (lang_hooks.decls.omp_report_decl (decl)), rkind);
-      inform (ctx->location, "enclosing OpenACC %qs construct", rkind);
-    }
-  else if (ctx->default_kind == OMP_CLAUSE_DEFAULT_PRESENT)
+	     DECL_NAME (lang_hooks.decls.omp_report_decl (decl)),
+	     oacc_region_type_name (ctx->region_type));
+      inform (ctx->oacc_default_clause_ctx->location,
+	      "enclosing OpenACC %qs construct",
+	      oacc_region_type_name
+	      (ctx->oacc_default_clause_ctx->region_type));
+    }
+  else if (default_kind == OMP_CLAUSE_DEFAULT_PRESENT)
     ; /* Handled above.  */
   else
-    gcc_checking_assert (ctx->default_kind == OMP_CLAUSE_DEFAULT_SHARED);
+    gcc_checking_assert (default_kind == OMP_CLAUSE_DEFAULT_SHARED);
 
   return flags;
 }
@@ -12124,6 +12152,10 @@  gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 
 	case OMP_CLAUSE_DEFAULT:
 	  ctx->default_kind = OMP_CLAUSE_DEFAULT_KIND (c);
+	  if (flag_openacc)
+	    /* For OpenACC, set current context to the 'active' one for
+	       default-clause lookup.  */
+	    ctx->oacc_default_clause_ctx = ctx;
 	  break;
 
 	case OMP_CLAUSE_INCLUSIVE:
diff --git a/gcc/testsuite/c-c++-common/goacc/default-3.c b/gcc/testsuite/c-c++-common/goacc/default-3.c
index ac169a903e9..e68f33606fd 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,49 @@  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) /* { dg-message "enclosing OpenACC .data. construct" } */
+#pragma acc kernels
+  {
+    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) /* { dg-message "enclosing OpenACC .data. construct" } */
+#pragma acc parallel
+  {
+    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 parallel default (none) /* { 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" } */
+  }
+
+#pragma acc data default (none) /* { dg-message "enclosing OpenACC .data. construct" } */
+#pragma acc data
+#pragma acc data
+#pragma acc parallel
+  {
+    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
+#pragma acc data default (none) /* { dg-message "enclosing OpenACC .data. construct" } */
+#pragma acc data
+#pragma acc parallel
+  {
+    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
+#pragma acc data
+#pragma acc data default (none) /* { dg-message "enclosing OpenACC .data. construct" } */
+#pragma acc parallel
+  {
+    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..76526396124 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 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 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..09b8c1bbd87 100644
--- a/gcc/testsuite/gfortran.dg/goacc/default-3.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/default-3.f95
@@ -15,4 +15,65 @@  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) ! { dg-message "enclosing OpenACC .data. construct" }
+  !$acc kernels
+  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) ! { dg-message "enclosing OpenACC .data. construct" }
+  !$acc parallel
+  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
+
+  !$acc data default (none)
+  !$acc parallel default (none) ! { 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
+
+  !$acc data default (none) ! { dg-message "enclosing OpenACC .data. construct" }
+  !$acc data
+  !$acc data
+  !$acc parallel
+  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
+  !$acc end data
+  !$acc end data
+
+  !$acc data
+  !$acc data default (none) ! { dg-message "enclosing OpenACC .data. construct" }
+  !$acc data
+  !$acc parallel
+  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
+  !$acc end data
+  !$acc end data
+
+  !$acc data
+  !$acc data
+  !$acc data default (none) ! { dg-message "enclosing OpenACC .data. construct" }
+  !$acc parallel
+  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
+  !$acc end data
+  !$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..9260db07a1a 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 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