Add OpenACC Fortran support for deviceptr and variable in common blocks

2018-06-29  Cesar Philippidis  <cesar@codesourcery.com>
	    James Norris  <jnorris@codesourcery.com>

	gcc/fortran/
	* openmp.cc (gfc_match_omp_map_clause): Re-write handling of the
	deviceptr clause.  Add new common_blocks argument.  Propagate it to
	gfc_match_omp_variable_list.
	(gfc_match_omp_clauses): Update calls to gfc_match_omp_map_clauses.
	(resolve_positive_int_expr): Promote the warning to an error.
	(check_array_not_assumed): Remove pointer check.
	(resolve_oacc_nested_loops): Error on do concurrent loops.
	* trans-openmp.cc (gfc_omp_finish_clause): Don't create pointer data
	mappings for deviceptr clauses.
	(gfc_trans_omp_clauses): Likewise.

	gcc/
	* gimplify.cc (enum gimplify_omp_var_data): Add GOVD_DEVICETPR.
	(oacc_default_clause): Privatize fortran common blocks.
	(omp_notice_variable): Add GOVD_DEVICEPTR attribute when appropriate.
	Defer the expansion of DECL_VALUE_EXPR for common block decls.
	(gimplify_scan_omp_clauses): Add GOVD_DEVICEPTR attribute when
	appropriate.
	(gimplify_adjust_omp_clauses_1): Set GOMP_MAP_FORCE_DEVICEPTR for
	implicit deviceptr mappings.

	gcc/testsuite/
	* c-c++-common/goacc/deviceptr-4.c: Update.
	* gfortran.dg/goacc/loop-2-kernels-tile.f95: Update.
	* gfortran.dg/goacc/loop-2-parallel-tile.f95: Update.
	* gfortran.dg/goacc/sie.f95: Update.
	* gfortran.dg/goacc/tile-1.f90: Update.
	* gfortran.dg/gomp/pr77516.f90: Update.

	libgomp/
	* oacc-parallel.c (GOACC_parallel_keyed): Handle Fortran deviceptr
	clause.
	(GOACC_data_start): Likewise.
	* testsuite/libgomp.oacc-fortran/deviceptr-1.f90: New test.

12 files changed, 80 insertions, 35 deletions

diff --git a/gcc/ChangeLog.omp b/gcc/ChangeLog.omp
index eed10a5..30b4ab9 100644
--- a/gcc/ChangeLog.omp
+++ b/gcc/ChangeLog.omp
@@ -1,3 +1,13 @@
+2018-06-29  Cesar Philippidis  <cesar@codesourcery.com>
+	    James Norris  <jnorris@codesourcery.com>
+
+	* gimplify.cc (enum gimplify_omp_var_data): Add GOVD_DEVICETPR.
+	(omp_notice_variable): Add GOVD_DEVICEPTR attribute when appropriate.
+	(gimplify_scan_omp_clauses): Add GOVD_DEVICEPTR attribute when
+	appropriate.
+	(gimplify_adjust_omp_clauses_1): Set GOMP_MAP_FORCE_DEVICEPTR for
+	implicit deviceptr mappings.
+
 2020-04-19  Chung-Lin Tang  <cltang@codesourcery.com>
 
 	PR other/76739
diff --git a/gcc/fortran/ChangeLog.omp b/gcc/fortran/ChangeLog.omp
index b742e1b..379b2ae 100644
--- a/gcc/fortran/ChangeLog.omp
+++ b/gcc/fortran/ChangeLog.omp
@@ -1,3 +1,9 @@
+2018-06-29  Cesar Philippidis  <cesar@codesourcery.com>
+	    James Norris  <jnorris@codesourcery.com>
+
+	* openmp.cc (resolve_positive_int_expr): Promote the warning to an
+	error.
+
 2020-04-19  Chung-Lin Tang  <cltang@codesourcery.com>
 
 	PR other/76739
diff --git a/gcc/fortran/openmp.cc b/gcc/fortran/openmp.cc
index 7141481..838fe2f 100644
--- a/gcc/fortran/openmp.cc
+++ b/gcc/fortran/openmp.cc
@@ -6071,8 +6071,8 @@ resolve_positive_int_expr (gfc_expr *expr, const char *clause)
   if (expr->expr_type == EXPR_CONSTANT
       && expr->ts.type == BT_INTEGER
       && mpz_sgn (expr->value.integer) <= 0)
-    gfc_warning (0, "INTEGER expression of %s clause at %L must be positive",
-		 clause, &expr->where);
+    gfc_error ("INTEGER expression of %s clause at %L must be positive",
+	       clause, &expr->where);
 }
 
 static void
diff --git a/gcc/fortran/trans-openmp.cc b/gcc/fortran/trans-openmp.cc
index 43d59ab..436694c 100644
--- a/gcc/fortran/trans-openmp.cc
+++ b/gcc/fortran/trans-openmp.cc
@@ -1501,6 +1501,9 @@ gfc_omp_finish_clause (tree c, gimple_seq *pre_p, bool openacc)
       return;
     }
 
+  if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FORCE_DEVICEPTR)
+    return;
+
   tree c2 = NULL_TREE, c3 = NULL_TREE, c4 = NULL_TREE;
   tree present = gfc_omp_check_optional_argument (decl, true);
   if (POINTER_TYPE_P (TREE_TYPE (decl)))
@@ -3104,6 +3107,12 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
 		      goto finalize_map_clause;
 		    }
 		  else if (POINTER_TYPE_P (TREE_TYPE (decl))
+			   && n->u.map_op == OMP_MAP_FORCE_DEVICEPTR)
+		    {
+		      OMP_CLAUSE_DECL (node) = decl;
+		      goto finalize_map_clause;
+		    }
+		  else if (POINTER_TYPE_P (TREE_TYPE (decl))
 			   && (gfc_omp_privatize_by_reference (decl)
 			       || GFC_DECL_GET_SCALAR_POINTER (decl)
 			       || GFC_DECL_GET_SCALAR_ALLOCATABLE (decl)
diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc
index 8daa552..e654031 100644
--- a/gcc/gimplify.cc
+++ b/gcc/gimplify.cc
@@ -132,6 +132,9 @@ enum gimplify_omp_var_data
   /* Flag for GOVD_FIRSTPRIVATE: OMP_CLAUSE_FIRSTPRIVATE_IMPLICIT.  */
   GOVD_FIRSTPRIVATE_IMPLICIT = 0x8000000,
 
+  /* Flag for OpenACC deviceptrs.  */
+  GOVD_DEVICEPTR = (1<<24),
+
   GOVD_DATA_SHARE_CLASS = (GOVD_SHARED | GOVD_PRIVATE | GOVD_FIRSTPRIVATE
 			   | GOVD_LASTPRIVATE | GOVD_REDUCTION | GOVD_LINEAR
 			   | GOVD_LOCAL)
@@ -7864,6 +7867,7 @@ omp_notice_variable (struct gimplify_omp_ctx *ctx, tree decl, bool in_code)
 		        error ("variable %qE declared in enclosing "
 			       "%<host_data%> region", DECL_NAME (decl));
 		      nflags |= GOVD_MAP;
+		      nflags |= (n2->value & GOVD_DEVICEPTR);
 		      if (octx->region_type == ORT_ACC_DATA
 			  && (n2->value & GOVD_MAP_0LEN_ARRAY))
 			nflags |= GOVD_MAP_0LEN_ARRAY;
@@ -10286,6 +10290,8 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 	  if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_TO
 	      || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_TOFROM)
 	    flags |= GOVD_MAP_ALWAYS_TO;
+	  else if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FORCE_DEVICEPTR)
+	    flags |= GOVD_DEVICEPTR;
 
 	  if ((code == OMP_TARGET
 	       || code == OMP_TARGET_DATA
@@ -11243,7 +11249,8 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n, void *data)
 		       | GOVD_MAP_FORCE
 		       | GOVD_MAP_FORCE_PRESENT
 		       | GOVD_MAP_ALLOC_ONLY
-		       | GOVD_MAP_FROM_ONLY))
+		       | GOVD_MAP_FROM_ONLY
+		       | GOVD_DEVICEPTR))
 	{
 	case 0:
 	  kind = GOMP_MAP_TOFROM;
@@ -11266,6 +11273,9 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n, void *data)
 	case GOVD_MAP_FORCE_PRESENT:
 	  kind = GOMP_MAP_FORCE_PRESENT;
 	  break;
+	case GOVD_DEVICEPTR:
+	  kind = GOMP_MAP_FORCE_DEVICEPTR;
+	  break;
 	default:
 	  gcc_unreachable ();
 	}
diff --git a/gcc/testsuite/ChangeLog.omp b/gcc/testsuite/ChangeLog.omp
index 64bb0cb..75d810f 100644
--- a/gcc/testsuite/ChangeLog.omp
+++ b/gcc/testsuite/ChangeLog.omp
@@ -1,3 +1,13 @@
+2018-06-29  Cesar Philippidis  <cesar@codesourcery.com>
+	    James Norris  <jnorris@codesourcery.com>
+
+	* c-c++-common/goacc/deviceptr-4.c: Update.
+	* gfortran.dg/goacc/loop-2-kernels-tile.f95: Update.
+	* gfortran.dg/goacc/loop-2-parallel-tile.f95: Update.
+	* gfortran.dg/goacc/sie.f95: Update.
+	* gfortran.dg/goacc/tile-1.f90: Update.
+	* gfortran.dg/gomp/pr77516.f90: Update.
+
 2020-04-19  Chung-Lin Tang  <cltang@codesourcery.com>
 
 	PR other/76739
diff --git a/gcc/testsuite/c-c++-common/goacc/deviceptr-4.c b/gcc/testsuite/c-c++-common/goacc/deviceptr-4.c
index db1b916..79a5162 100644
--- a/gcc/testsuite/c-c++-common/goacc/deviceptr-4.c
+++ b/gcc/testsuite/c-c++-common/goacc/deviceptr-4.c
@@ -8,4 +8,4 @@ subr (int *a)
   a[0] += 1.0;
 }
 
-/* { dg-final { scan-tree-dump-times "#pragma omp target oacc_parallel.*map\\(tofrom:a" 1 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "#pragma omp target oacc_parallel.*map\\(force_deviceptr:a" 1 "gimple" } } */
diff --git a/gcc/testsuite/gfortran.dg/goacc/loop-2-kernels-tile.f95 b/gcc/testsuite/gfortran.dg/goacc/loop-2-kernels-tile.f95
index afc8a27..6542515 100644
--- a/gcc/testsuite/gfortran.dg/goacc/loop-2-kernels-tile.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/loop-2-kernels-tile.f95
@@ -29,7 +29,7 @@ program test
       DO j = 1,10
       ENDDO
     ENDDO
-    !$acc loop tile(-1) ! { dg-warning "must be positive" }
+    !$acc loop tile(-1) ! { dg-error "must be positive" }
     do i = 1,10
     enddo
     !$acc loop tile(i) ! { dg-error "constant expression" }
@@ -82,7 +82,7 @@ program test
     DO j = 1,10
     ENDDO
   ENDDO
-  !$acc kernels loop tile(-1) ! { dg-warning "must be positive" }
+  !$acc kernels loop tile(-1) ! { dg-error "must be positive" }
   do i = 1,10
   enddo
   !$acc kernels loop tile(i) ! { dg-error "constant expression" }
diff --git a/gcc/testsuite/gfortran.dg/goacc/loop-2-parallel-tile.f95 b/gcc/testsuite/gfortran.dg/goacc/loop-2-parallel-tile.f95
index 4bfca74..dae8f66 100644
--- a/gcc/testsuite/gfortran.dg/goacc/loop-2-parallel-tile.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/loop-2-parallel-tile.f95
@@ -20,7 +20,7 @@ program test
       DO j = 1,10
       ENDDO
     ENDDO
-    !$acc loop tile(-1) ! { dg-warning "must be positive" }
+    !$acc loop tile(-1) ! { dg-error "must be positive" }
     do i = 1,10
     enddo
     !$acc loop tile(i) ! { dg-error "constant expression" }
@@ -73,7 +73,7 @@ program test
     DO j = 1,10
     ENDDO
   ENDDO
-  !$acc parallel loop tile(-1) ! { dg-warning "must be positive" }
+  !$acc parallel loop tile(-1) ! { dg-error "must be positive" }
   do i = 1,10
   enddo
   !$acc parallel loop tile(i) ! { dg-error "constant expression" }
diff --git a/gcc/testsuite/gfortran.dg/goacc/sie.f95 b/gcc/testsuite/gfortran.dg/goacc/sie.f95
index 5982d5d..f393cf2 100644
--- a/gcc/testsuite/gfortran.dg/goacc/sie.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/sie.f95
@@ -78,10 +78,10 @@ program test
   !$acc parallel num_gangs(i+1)
   !$acc end parallel
 
-  !$acc parallel num_gangs(-1) ! { dg-warning "must be positive" }
+  !$acc parallel num_gangs(-1) ! { dg-error "must be positive" }
   !$acc end parallel
 
-  !$acc parallel num_gangs(0) ! { dg-warning "must be positive" }
+  !$acc parallel num_gangs(0) ! { dg-error "must be positive" }
   !$acc end parallel
 
   !$acc parallel num_gangs() ! { dg-error "Invalid character in name" }
@@ -106,10 +106,10 @@ program test
   !$acc kernels num_gangs(i+1)
   !$acc end kernels
 
-  !$acc kernels num_gangs(-1) ! { dg-warning "must be positive" }
+  !$acc kernels num_gangs(-1) ! { dg-error "must be positive" }
   !$acc end kernels
 
-  !$acc kernels num_gangs(0) ! { dg-warning "must be positive" }
+  !$acc kernels num_gangs(0) ! { dg-error "must be positive" }
   !$acc end kernels
 
   !$acc kernels num_gangs() ! { dg-error "Invalid character in name" }
@@ -135,10 +135,10 @@ program test
   !$acc parallel num_workers(i+1)
   !$acc end parallel
 
-  !$acc parallel num_workers(-1) ! { dg-warning "must be positive" }
+  !$acc parallel num_workers(-1) ! { dg-error "must be positive" }
   !$acc end parallel
 
-  !$acc parallel num_workers(0) ! { dg-warning "must be positive" }
+  !$acc parallel num_workers(0) ! { dg-error "must be positive" }
   !$acc end parallel
 
   !$acc parallel num_workers() ! { dg-error "Invalid expression after 'num_workers\\('" }
@@ -163,10 +163,10 @@ program test
   !$acc kernels num_workers(i+1)
   !$acc end kernels
 
-  !$acc kernels num_workers(-1) ! { dg-warning "must be positive" }
+  !$acc kernels num_workers(-1) ! { dg-error "must be positive" }
   !$acc end kernels
 
-  !$acc kernels num_workers(0) ! { dg-warning "must be positive" }
+  !$acc kernels num_workers(0) ! { dg-error "must be positive" }
   !$acc end kernels
 
   !$acc kernels num_workers() ! { dg-error "Invalid expression after 'num_workers\\('" }
@@ -192,10 +192,10 @@ program test
   !$acc parallel vector_length(i+1)
   !$acc end parallel
 
-  !$acc parallel vector_length(-1) ! { dg-warning "must be positive" }
+  !$acc parallel vector_length(-1) ! { dg-error "must be positive" }
   !$acc end parallel
 
-  !$acc parallel vector_length(0) ! { dg-warning "must be positive" }
+  !$acc parallel vector_length(0) ! { dg-error "must be positive" }
   !$acc end parallel
 
   !$acc parallel vector_length() ! { dg-error "Invalid expression after 'vector_length\\('" }
@@ -220,10 +220,10 @@ program test
   !$acc kernels vector_length(i+1)
   !$acc end kernels
 
-  !$acc kernels vector_length(-1) ! { dg-warning "must be positive" }
+  !$acc kernels vector_length(-1) ! { dg-error "must be positive" }
   !$acc end kernels
 
-  !$acc kernels vector_length(0) ! { dg-warning "must be positive" }
+  !$acc kernels vector_length(0) ! { dg-error "must be positive" }
   !$acc end kernels
 
   !$acc kernels vector_length() ! { dg-error "Invalid expression after 'vector_length\\('" }
@@ -250,10 +250,10 @@ program test
   !$acc loop gang(i+1)
   do i = 1,10
   enddo
-  !$acc loop gang(-1) ! { dg-warning "must be positive" }
+  !$acc loop gang(-1) ! { dg-error "must be positive" }
   do i = 1,10
   enddo
-  !$acc loop gang(0) ! { dg-warning "must be positive" }
+  !$acc loop gang(0) ! { dg-error "must be positive" }
   do i = 1,10
   enddo
   !$acc loop gang() ! { dg-error "Invalid character in name" }
@@ -282,10 +282,10 @@ program test
   !$acc loop worker(i+1)
   do i = 1,10
   enddo
-  !$acc loop worker(-1) ! { dg-warning "must be positive" }
+  !$acc loop worker(-1) ! { dg-error "must be positive" }
   do i = 1,10
   enddo
-  !$acc loop worker(0) ! { dg-warning "must be positive" }
+  !$acc loop worker(0) ! { dg-error "must be positive" }
   do i = 1,10
   enddo
   !$acc loop worker() ! { dg-error "Invalid character in name" }
@@ -314,10 +314,10 @@ program test
   !$acc loop vector(i+1)
   do i = 1,10
   enddo
-  !$acc loop vector(-1) ! { dg-warning "must be positive" }
+  !$acc loop vector(-1) ! { dg-error "must be positive" }
   do i = 1,10
   enddo
-  !$acc loop vector(0) ! { dg-warning "must be positive" }
+  !$acc loop vector(0) ! { dg-error "must be positive" }
   do i = 1,10
   enddo
   !$acc loop vector() ! { dg-error "Invalid character in name" }
diff --git a/gcc/testsuite/gfortran.dg/goacc/tile-1.f90 b/gcc/testsuite/gfortran.dg/goacc/tile-1.f90
index f609b12..9ef7521 100644
--- a/gcc/testsuite/gfortran.dg/goacc/tile-1.f90
+++ b/gcc/testsuite/gfortran.dg/goacc/tile-1.f90
@@ -44,17 +44,17 @@ subroutine parloop
   do i = 1, n
   end do
 
-  !$acc parallel loop tile(-3) ! { dg-warning "must be positive" }
+  !$acc parallel loop tile(-3) ! { dg-error "must be positive" }
   do i = 1, n
   end do
 
-  !$acc parallel loop tile(10, -3) ! { dg-warning "must be positive" }
+  !$acc parallel loop tile(10, -3) ! { dg-error "must be positive" }
   do i = 1, n
      do j = 1, n
      end do
   end do
 
-  !$acc parallel loop tile(-100, 10, 5) ! { dg-warning "must be positive" }
+  !$acc parallel loop tile(-100, 10, 5) ! { dg-error "must be positive" }
   do i = 1, n
      do j = 1, n
         do k = 1, n
@@ -114,7 +114,7 @@ subroutine par
      end do
   end do
 
-  !$acc loop tile(-2) ! { dg-warning "must be positive" }
+  !$acc loop tile(-2) ! { dg-error "must be positive" }
   do i = 1, n
   end do
 
@@ -195,7 +195,7 @@ subroutine kern
      end do
   end do
 
-  !$acc loop tile(-2) ! { dg-warning "must be positive" }
+  !$acc loop tile(-2) ! { dg-error "must be positive" }
   do i = 1, n
   end do
 
@@ -295,17 +295,17 @@ subroutine kernsloop
   do i = 1, n
   end do
 
-  !$acc kernels loop tile(-3) ! { dg-warning "must be positive" }
+  !$acc kernels loop tile(-3) ! { dg-error "must be positive" }
   do i = 1, n
   end do
 
-  !$acc kernels loop tile(10, -3) ! { dg-warning "must be positive" }
+  !$acc kernels loop tile(10, -3) ! { dg-error "must be positive" }
   do i = 1, n
      do j = 1, n
      end do
   end do
 
-  !$acc kernels loop tile(-100, 10, 5) ! { dg-warning "must be positive" }
+  !$acc kernels loop tile(-100, 10, 5) ! { dg-error "must be positive" }
   do i = 1, n
      do j = 1, n
         do k = 1, n
diff --git a/gcc/testsuite/gfortran.dg/gomp/pr77516.f90 b/gcc/testsuite/gfortran.dg/gomp/pr77516.f90
index 9c0a95b..3ac3f55 100644
--- a/gcc/testsuite/gfortran.dg/gomp/pr77516.f90
+++ b/gcc/testsuite/gfortran.dg/gomp/pr77516.f90
@@ -4,7 +4,7 @@
 program pr77516
    integer :: i, x
    x = 0
-!$omp simd safelen(0) reduction(+:x)	! { dg-warning "must be positive" }
+!$omp simd safelen(0) reduction(+:x)	! { dg-error "must be positive" }
    do i = 1, 8
       x = x + 1
    end do