Re: [patch] Add OpenACC Fortran support for deviceptr and variable in common blocks

[Chung-Lin Tang <chunglin_tang@mentor.com>]

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

Hi Jakub,
I have updated this patch and rebased with current trunk. The issues you identified
in the last review should all be resolved. Is this now okay for trunk?

Thanks, and Happy New Year~
Chung-Lin


[-- Attachment #2: openacc-fortran-deviceptr-commonblk.v2.patch --]
[-- Type: text/plain, Size: 39083 bytes --]

Index: gcc/fortran/openmp.c
===================================================================
--- gcc/fortran/openmp.c	(revision 267513)
+++ gcc/fortran/openmp.c	(working copy)
@@ -914,10 +914,11 @@ omp_inv_mask::omp_inv_mask (const omp_mask &m) : o
    mapping.  */
 
 static bool
-gfc_match_omp_map_clause (gfc_omp_namelist **list, gfc_omp_map_op map_op)
+gfc_match_omp_map_clause (gfc_omp_namelist **list, gfc_omp_map_op map_op,
+			  bool common_blocks = true)
 {
   gfc_omp_namelist **head = NULL;
-  if (gfc_match_omp_variable_list ("", list, false, NULL, &head, true)
+  if (gfc_match_omp_variable_list ("", list, common_blocks, NULL, &head, true)
       == MATCH_YES)
     {
       gfc_omp_namelist *n;
@@ -1156,12 +1157,12 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const
 	      && openacc
 	      && gfc_match ("device ( ") == MATCH_YES
 	      && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
-					   OMP_MAP_FORCE_TO))
+					   OMP_MAP_FORCE_TO, false))
 	    continue;
 	  if ((mask & OMP_CLAUSE_DEVICEPTR)
 	      && gfc_match ("deviceptr ( ") == MATCH_YES
 	      && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
-					   OMP_MAP_FORCE_DEVICEPTR))
+					   OMP_MAP_FORCE_DEVICEPTR, false))
 	    continue;
 	  if ((mask & OMP_CLAUSE_DEVICE_RESIDENT)
 	      && gfc_match_omp_variable_list
@@ -1531,7 +1532,7 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const
 	  if ((mask & OMP_CLAUSE_PRESENT)
 	      && gfc_match ("present ( ") == MATCH_YES
 	      && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
-					   OMP_MAP_FORCE_PRESENT))
+					   OMP_MAP_FORCE_PRESENT, false))
 	    continue;
 	  if ((mask & OMP_CLAUSE_COPY)
 	      && gfc_match ("present_or_copy ( ") == MATCH_YES
@@ -3762,7 +3763,7 @@ check_symbol_not_pointer (gfc_symbol *sym, locus l
 /* Emits error when symbol represents assumed size/rank array.  */
 
 static void
-check_array_not_assumed (gfc_symbol *sym, locus loc, const char *name)
+check_oacc_array_not_assumed (gfc_symbol *sym, locus loc, const char *name)
 {
   if (sym->as && sym->as->type == AS_ASSUMED_SIZE)
     gfc_error ("Assumed size array %qs in %s clause at %L",
@@ -3770,10 +3771,6 @@ static void
   if (sym->as && sym->as->type == AS_ASSUMED_RANK)
     gfc_error ("Assumed rank array %qs in %s clause at %L",
 	       sym->name, name, &loc);
-  if (sym->as && sym->as->type == AS_DEFERRED && sym->attr.pointer
-      && !sym->attr.contiguous)
-    gfc_error ("Noncontiguous deferred shape array %qs in %s clause at %L",
-	       sym->name, name, &loc);
 }
 
 static void
@@ -3788,7 +3785,7 @@ resolve_oacc_data_clauses (gfc_symbol *sym, locus
     gfc_error ("ALLOCATABLE object %qs of polymorphic type "
 	       "in %s clause at %L", sym->name, name, &loc);
   check_symbol_not_pointer (sym, loc, name);
-  check_array_not_assumed (sym, loc, name);
+  check_oacc_array_not_assumed (sym, loc, name);
 }
 
 static void
@@ -3817,7 +3814,7 @@ resolve_oacc_deviceptr_clause (gfc_symbol *sym, lo
   if (sym->attr.value)
     gfc_error ("VALUE object %qs in %s clause at %L",
 	       sym->name, name, &loc);
-  check_array_not_assumed (sym, loc, name);
+  check_oacc_array_not_assumed (sym, loc, name);
 }
 
 
@@ -4508,7 +4505,7 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_claus
 		  }
 		if (code
 		    && (oacc_is_loop (code) || code->op == EXEC_OACC_PARALLEL))
-		  check_array_not_assumed (n->sym, n->where, name);
+		  check_oacc_array_not_assumed (n->sym, n->where, name);
 		else if (n->sym->as && n->sym->as->type == AS_ASSUMED_SIZE)
 		  gfc_error ("Assumed size array %qs in %s clause at %L",
 			     n->sym->name, name, &n->where);
@@ -4725,7 +4722,7 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_claus
 		      /* FALLTHRU */
 		  case OMP_LIST_DEVICE_RESIDENT:
 		    check_symbol_not_pointer (n->sym, n->where, name);
-		    check_array_not_assumed (n->sym, n->where, name);
+		    check_oacc_array_not_assumed (n->sym, n->where, name);
 		    break;
 		  default:
 		    break;
@@ -5760,7 +5757,13 @@ resolve_oacc_nested_loops (gfc_code *code, gfc_cod
 		     "at %L", &do_code->loc);
 	  break;
 	}
-      gcc_assert (do_code->op == EXEC_DO || do_code->op == EXEC_DO_CONCURRENT);
+      if (do_code->op == EXEC_DO_CONCURRENT)
+	{
+	  gfc_error ("!$ACC LOOP cannot be a DO CONCURRENT loop at %L",
+		     &do_code->loc);
+	  break;
+	}
+      gcc_assert (do_code->op == EXEC_DO);
       if (do_code->ext.iterator->var->ts.type != BT_INTEGER)
 	gfc_error ("!$ACC LOOP iteration variable must be of type integer at %L",
 		   &do_code->loc);
@@ -5878,7 +5881,12 @@ resolve_oacc_loop_blocks (gfc_code *code)
 	    }
 	  else
 	    {
-	      resolve_positive_int_expr (el->expr, "TILE");
+	      resolve_scalar_int_expr (el->expr, "TILE");
+	      if (el->expr->expr_type == EXPR_CONSTANT
+		  && el->expr->ts.type == BT_INTEGER
+		  && mpz_sgn (el->expr->value.integer) <= 0)
+		gfc_error ("INTEGER expression of TILE clause at %L must be "
+			   "positive", &el->expr->where);
 	      if (el->expr->expr_type != EXPR_CONSTANT)
 		gfc_error ("TILE requires constant expression at %L",
 			   &code->loc);
@@ -5965,7 +5973,7 @@ gfc_resolve_oacc_declare (gfc_namespace *ns)
 	  }
 
       for (n = oc->clauses->lists[OMP_LIST_DEVICE_RESIDENT]; n; n = n->next)
-	check_array_not_assumed (n->sym, oc->loc, "DEVICE_RESIDENT");
+	check_oacc_array_not_assumed (n->sym, oc->loc, "DEVICE_RESIDENT");
     }
 
   for (oc = ns->oacc_declare; oc; oc = oc->next)
Index: gcc/fortran/trans-openmp.c
===================================================================
--- gcc/fortran/trans-openmp.c	(revision 267513)
+++ gcc/fortran/trans-openmp.c	(working copy)
@@ -1066,6 +1066,8 @@ gfc_omp_finish_clause (tree c, gimple_seq *pre_p)
     }
 
   tree c2 = NULL_TREE, c3 = NULL_TREE, c4 = NULL_TREE;
+  if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FORCE_DEVICEPTR)
+    return;
   if (POINTER_TYPE_P (TREE_TYPE (decl)))
     {
       if (!gfc_omp_privatize_by_reference (decl)
@@ -2117,6 +2119,12 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp
 	      if (n->expr == NULL || n->expr->ref->u.ar.type == AR_FULL)
 		{
 		  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)
@@ -2288,6 +2296,7 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp
 		  ptr2 = fold_convert (sizetype, ptr2);
 		  OMP_CLAUSE_SIZE (node3)
 		    = fold_build2 (MINUS_EXPR, sizetype, ptr, ptr2);
+		finalize_map_clause:;
 		}
 	      switch (n->u.map_op)
 		{
Index: gcc/gimplify.c
===================================================================
--- gcc/gimplify.c	(revision 267513)
+++ gcc/gimplify.c	(working copy)
@@ -113,6 +113,9 @@ enum gimplify_omp_var_data
 
   GOVD_NONTEMPORAL = 4194304,
 
+  /* Flag for OpenACC deviceptrs.  */
+  GOVD_DEVICEPTR = 8388608,
+
   GOVD_DATA_SHARE_CLASS = (GOVD_SHARED | GOVD_PRIVATE | GOVD_FIRSTPRIVATE
 			   | GOVD_LASTPRIVATE | GOVD_REDUCTION | GOVD_LINEAR
 			   | GOVD_LOCAL)
@@ -7109,6 +7112,7 @@ oacc_default_clause (struct gimplify_omp_ctx *ctx,
 {
   const char *rkind;
   bool on_device = false;
+  bool is_private = false;
   bool declared = is_oacc_declared (decl);
   tree type = TREE_TYPE (decl);
 
@@ -7115,9 +7119,13 @@ oacc_default_clause (struct gimplify_omp_ctx *ctx,
   if (lang_hooks.decls.omp_privatize_by_reference (decl))
     type = TREE_TYPE (type);
 
+  if (RECORD_OR_UNION_TYPE_P (type))
+    is_private = lang_hooks.decls.omp_disregard_value_expr (decl, false);
+
   if ((ctx->region_type & (ORT_ACC_PARALLEL | ORT_ACC_KERNELS)) != 0
       && is_global_var (decl)
-      && device_resident_p (decl))
+      && device_resident_p (decl)
+      && !is_private)
     {
       on_device = true;
       flags |= GOVD_MAP_TO_ONLY;
@@ -7128,7 +7136,9 @@ oacc_default_clause (struct gimplify_omp_ctx *ctx,
     case ORT_ACC_KERNELS:
       rkind = "kernels";
 
-      if (AGGREGATE_TYPE_P (type))
+      if (is_private)
+	flags |= GOVD_MAP;
+      else if (AGGREGATE_TYPE_P (type))
 	{
 	  /* Aggregates default to 'present_or_copy', or 'present'.  */
 	  if (ctx->default_kind != OMP_CLAUSE_DEFAULT_PRESENT)
@@ -7145,7 +7155,9 @@ oacc_default_clause (struct gimplify_omp_ctx *ctx,
     case ORT_ACC_PARALLEL:
       rkind = "parallel";
 
-      if (on_device || declared)
+      if (is_private)
+	flags |= GOVD_FIRSTPRIVATE;
+      else if (on_device || declared)
 	flags |= GOVD_MAP;
       else if (AGGREGATE_TYPE_P (type))
 	{
@@ -7211,7 +7223,8 @@ omp_notice_variable (struct gimplify_omp_ctx *ctx,
 	{
 	  tree value = get_base_address (DECL_VALUE_EXPR (decl));
 
-	  if (value && DECL_P (value) && DECL_THREAD_LOCAL_P (value))
+	  if (!(ctx->region_type & ORT_ACC)
+	      && value && DECL_P (value) && DECL_THREAD_LOCAL_P (value))
 	    return omp_notice_threadprivate_variable (ctx, decl, value);
 	}
 
@@ -7243,7 +7256,8 @@ omp_notice_variable (struct gimplify_omp_ctx *ctx,
   n = splay_tree_lookup (ctx->variables, (splay_tree_key)decl);
   if ((ctx->region_type & ORT_TARGET) != 0)
     {
-      ret = lang_hooks.decls.omp_disregard_value_expr (decl, true);
+      shared = !(ctx->region_type & ORT_ACC);
+      ret = lang_hooks.decls.omp_disregard_value_expr (decl, shared);
       if (n == NULL)
 	{
 	  unsigned nflags = flags;
@@ -7316,6 +7330,7 @@ omp_notice_variable (struct gimplify_omp_ctx *ctx,
 		        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;
@@ -7411,6 +7426,8 @@ omp_notice_variable (struct gimplify_omp_ctx *ctx,
     }
 
   shared = ((flags | n->value) & GOVD_SHARED) != 0;
+  if (ctx->region_type & ORT_ACC)
+    shared = false;
   ret = lang_hooks.decls.omp_disregard_value_expr (decl, shared);
 
   /* If nothing changed, there's nothing left to do.  */
@@ -8804,6 +8821,8 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_se
 	  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;
 	  goto do_add;
 
 	case OMP_CLAUSE_DEPEND:
@@ -9534,7 +9553,8 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n,
 		       | 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;
@@ -9557,6 +9577,9 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n,
 	case GOVD_MAP_FORCE_PRESENT:
 	  kind = GOMP_MAP_FORCE_PRESENT;
 	  break;
+	case GOVD_DEVICEPTR:
+	  kind = GOMP_MAP_FORCE_DEVICEPTR;
+	  break;
 	default:
 	  gcc_unreachable ();
 	}
Index: gcc/testsuite/c-c++-common/goacc/deviceptr-4.c
===================================================================
--- gcc/testsuite/c-c++-common/goacc/deviceptr-4.c	(revision 267513)
+++ gcc/testsuite/c-c++-common/goacc/deviceptr-4.c	(working copy)
@@ -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" } } */
Index: gcc/testsuite/gfortran.dg/goacc/common-block-1.f90
===================================================================
--- gcc/testsuite/gfortran.dg/goacc/common-block-1.f90	(nonexistent)
+++ gcc/testsuite/gfortran.dg/goacc/common-block-1.f90	(working copy)
@@ -0,0 +1,69 @@
+! Test data clauses involving common blocks and common block data.
+! Specifically, validate early matching errors.
+
+subroutine subtest
+  implicit none
+  integer, parameter :: n = 10
+  integer a(n), b(n), c, d(n), e
+  real*4 x(n), y(n), z, w(n), v
+  common /blockA/ a, c, x
+  common /blockB/ b, y, z
+  !$acc declare link(/blockA/, /blockB/, e, v)
+end subroutine subtest
+
+program test
+  implicit none
+  integer, parameter :: n = 10
+  integer a(n), b(n), c, d(n), e
+  real*4 x(n), y(n), z, w(n), v
+  common /blockA/ a, c, x
+  common /blockB/ b, y, z
+  !$acc declare link(/blockA/, /blockB/, e, v)
+
+  !$acc data copy(/blockA/, /blockB/, e, v)
+  !$acc end data
+
+  !$acc data copyin(/blockA/, /blockB/, e, v)
+  !$acc end data
+
+  !$acc data copyout(/blockA/, /blockB/, e, v)
+  !$acc end data
+
+  !$acc data create(/blockA/, /blockB/, e, v)
+  !$acc end data
+
+  !$acc data copyout(/blockA/, /blockB/, e, v)
+  !$acc end data
+
+  !$acc data pcopy(/blockA/, /blockB/, e, v)
+  !$acc end data
+
+  !$acc data pcopyin(/blockA/, /blockB/, e, v)
+  !$acc end data
+
+  !$acc data pcopyout(/blockA/, /blockB/, e, v)
+  !$acc end data
+
+  !$acc data pcreate(/blockA/, /blockB/, e, v)
+  !$acc end data
+
+  !$acc data pcopyout(/blockA/, /blockB/, e, v)
+  !$acc end data
+
+  !$acc data present(/blockA/, /blockB/, e, v) ! { dg-error "" }
+  !$acc end data ! { dg-error "Unexpected ..ACC END DATA statement" }
+
+  !$acc parallel private(/blockA/, /blockB/, e, v)
+  !$acc end parallel
+
+  !$acc parallel firstprivate(/blockA/, /blockB/, e, v)
+  !$acc end parallel
+
+  !$acc exit data delete(/blockA/, /blockB/, e, v)
+
+  !$acc data deviceptr(/blockA/, /blockB/, e, v) ! { dg-error "Syntax error" }
+  !$acc end data ! { dg-error "Unexpected ..ACC END DATA statement" }
+
+  !$acc data deviceptr(/blockA/, /blockB/, e, v) ! { dg-error "Syntax error" }
+  !$acc end data ! { dg-error "Unexpected ..ACC END DATA statement" }
+end program test
Index: gcc/testsuite/gfortran.dg/goacc/common-block-2.f90
===================================================================
--- gcc/testsuite/gfortran.dg/goacc/common-block-2.f90	(nonexistent)
+++ gcc/testsuite/gfortran.dg/goacc/common-block-2.f90	(working copy)
@@ -0,0 +1,49 @@
+! Test data clauses involving common blocks and common block data.
+! Specifically, resolver errors such as duplicate data clauses.
+
+program test
+  implicit none
+  integer, parameter :: n = 10
+  integer a(n), b(n), c, d(n), e
+  real*4 x(n), y(n), z, w(n), v
+  common /blockA/ a, c, x
+  common /blockB/ b, y, z
+
+  !$acc data copy(/blockA/, /blockB/, e, v, a) ! { dg-error "Symbol .a. present on multiple clauses" }
+  !$acc end data
+
+  !$acc data copyin(/blockA/, /blockB/, e, v, a) ! { dg-error "Symbol .a. present on multiple clauses" }
+  !$acc end data
+
+  !$acc data copyout(/blockA/, /blockB/, e, v, a) ! { dg-error "Symbol .a. present on multiple clauses" }
+  !$acc end data
+
+  !$acc data create(/blockA/, /blockB/, e, v, a) ! { dg-error "Symbol .a. present on multiple clauses" }
+  !$acc end data
+
+  !$acc data copyout(/blockA/, /blockB/, e, v, a) ! { dg-error "Symbol .a. present on multiple clauses" }
+  !$acc end data
+
+  !$acc data pcopy(/blockA/, /blockB/, e, v, a) ! { dg-error "Symbol .a. present on multiple clauses" }
+  !$acc end data
+
+  !$acc data pcopyin(/blockA/, /blockB/, e, v, a) ! { dg-error "Symbol .a. present on multiple clauses" }
+  !$acc end data
+
+  !$acc data pcopyout(/blockA/, /blockB/, e, v, a) ! { dg-error "Symbol .a. present on multiple clauses" }
+  !$acc end data
+
+  !$acc data pcreate(/blockA/, /blockB/, e, v, a) ! { dg-error "Symbol .a. present on multiple clauses" }
+  !$acc end data
+
+  !$acc data pcopyout(/blockA/, /blockB/, e, v, a) ! { dg-error "Symbol .a. present on multiple clauses" }
+  !$acc end data
+
+  !$acc parallel private(/blockA/, /blockB/, e, v, a) ! { dg-error "Symbol .a. present on multiple clauses" }
+  !$acc end parallel
+
+  !$acc parallel firstprivate(/blockA/, /blockB/, e, v, a) ! { dg-error "Symbol .a. present on multiple clauses" }
+  !$acc end parallel
+
+  !$acc exit data delete(/blockA/, /blockB/, e, v, a) ! { dg-error "Symbol .a. present on multiple clauses" }
+end program test
Index: gcc/testsuite/gfortran.dg/goacc/loop-2-kernels-tile.f95
===================================================================
--- gcc/testsuite/gfortran.dg/goacc/loop-2-kernels-tile.f95	(revision 267513)
+++ gcc/testsuite/gfortran.dg/goacc/loop-2-kernels-tile.f95	(working copy)
@@ -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" }
Index: gcc/testsuite/gfortran.dg/goacc/loop-2-parallel-tile.f95
===================================================================
--- gcc/testsuite/gfortran.dg/goacc/loop-2-parallel-tile.f95	(revision 267513)
+++ gcc/testsuite/gfortran.dg/goacc/loop-2-parallel-tile.f95	(working copy)
@@ -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" }
Index: gcc/testsuite/gfortran.dg/goacc/loop-3-2.f95
===================================================================
--- gcc/testsuite/gfortran.dg/goacc/loop-3-2.f95	(revision 267513)
+++ gcc/testsuite/gfortran.dg/goacc/loop-3-2.f95	(working copy)
@@ -27,9 +27,9 @@ subroutine test1
   !$acc end parallel
   !$acc end loop ! { dg-error "Unexpected" }
 
-  ! OpenACC supports Fortran 2008 do concurrent statement
+  ! OpenACC does not support Fortran 2008 do concurrent statement
   !$acc loop
-  do concurrent (i = 1:5)
+  do concurrent (i = 1:5) ! { dg-error "ACC LOOP cannot be a DO CONCURRENT loop" }
   end do
 
   !$acc loop
Index: gcc/testsuite/gfortran.dg/goacc/loop-3.f95
===================================================================
--- gcc/testsuite/gfortran.dg/goacc/loop-3.f95	(revision 267513)
+++ gcc/testsuite/gfortran.dg/goacc/loop-3.f95	(working copy)
@@ -24,9 +24,9 @@ subroutine test1
   !$acc end parallel
   !$acc end loop ! { dg-error "Unexpected" }
 
-  ! OpenACC supports Fortran 2008 do concurrent statement
+  ! OpenACC does not support Fortran 2008 do concurrent statement
   !$acc loop
-  do concurrent (i = 1:5)
+  do concurrent (i = 1:5) ! { dg-error "ACC LOOP cannot be a DO CONCURRENT loop" }
   end do
 
   !$acc loop
Index: gcc/testsuite/gfortran.dg/goacc/loop-5.f95
===================================================================
--- gcc/testsuite/gfortran.dg/goacc/loop-5.f95	(revision 267513)
+++ gcc/testsuite/gfortran.dg/goacc/loop-5.f95	(working copy)
@@ -93,9 +93,6 @@ program test
       DO j = 1,10
       ENDDO
     ENDDO
-    !$acc loop tile(-1) ! { dg-warning "must be positive" }
-    do i = 1,10
-    enddo
     !$acc loop vector tile(*)
     DO i = 1,10
     ENDDO
@@ -129,9 +126,6 @@ program test
       DO j = 1,10
       ENDDO
     ENDDO
-    !$acc loop tile(-1) ! { dg-warning "must be positive" }
-    do i = 1,10
-    enddo
     !$acc loop vector tile(*)
     DO i = 1,10
     ENDDO
@@ -242,9 +236,6 @@ program test
     DO j = 1,10
     ENDDO
   ENDDO
-  !$acc kernels loop tile(-1) ! { dg-warning "must be positive" }
-  do i = 1,10
-  enddo
   !$acc kernels loop vector tile(*)
   DO i = 1,10
   ENDDO
@@ -333,9 +324,6 @@ program test
     DO j = 1,10
     ENDDO
   ENDDO
-  !$acc parallel loop tile(-1) ! { dg-warning "must be positive" }
-  do i = 1,10
-  enddo
   !$acc parallel loop vector tile(*)
   DO i = 1,10
   ENDDO
Index: gcc/testsuite/gfortran.dg/goacc/pr72715.f90
===================================================================
--- gcc/testsuite/gfortran.dg/goacc/pr72715.f90	(nonexistent)
+++ gcc/testsuite/gfortran.dg/goacc/pr72715.f90	(working copy)
@@ -0,0 +1,6 @@
+program p
+  integer :: i
+  !$acc loop
+  do concurrent (i=1:3) ! { dg-error "ACC LOOP cannot be a DO CONCURRENT loop" }
+  end do
+end program p
Index: gcc/testsuite/gfortran.dg/goacc/sie.f95
===================================================================
--- gcc/testsuite/gfortran.dg/goacc/sie.f95	(revision 267513)
+++ gcc/testsuite/gfortran.dg/goacc/sie.f95	(working copy)
@@ -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 character in name" }
@@ -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 character in name" }
@@ -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 character in name" }
@@ -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 character in name" }
@@ -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" }
Index: gcc/testsuite/gfortran.dg/goacc/tile-1.f90
===================================================================
--- gcc/testsuite/gfortran.dg/goacc/tile-1.f90	(revision 267513)
+++ gcc/testsuite/gfortran.dg/goacc/tile-1.f90	(working copy)
@@ -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
Index: gcc/testsuite/gfortran.dg/gomp/pr77516.f90
===================================================================
--- gcc/testsuite/gfortran.dg/gomp/pr77516.f90	(revision 267513)
+++ gcc/testsuite/gfortran.dg/gomp/pr77516.f90	(working copy)
@@ -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
Index: libgomp/oacc-parallel.c
===================================================================
--- libgomp/oacc-parallel.c	(revision 267513)
+++ libgomp/oacc-parallel.c	(working copy)
@@ -312,6 +312,8 @@ GOACC_data_start (int flags_m, size_t mapnum,
   struct goacc_thread *thr = goacc_thread ();
   struct gomp_device_descr *acc_dev = thr->dev;
 
+  handle_ftn_pointers (mapnum, hostaddrs, sizes, kinds);
+
   /* Host fallback or 'do nothing'.  */
   if ((acc_dev->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM)
       || (flags & GOACC_FLAG_HOST_FALLBACK))
Index: libgomp/testsuite/libgomp.oacc-fortran/common-block-1.f90
===================================================================
--- libgomp/testsuite/libgomp.oacc-fortran/common-block-1.f90	(nonexistent)
+++ libgomp/testsuite/libgomp.oacc-fortran/common-block-1.f90	(working copy)
@@ -0,0 +1,105 @@
+! Test data located inside common blocks.  This test does not execrise
+! ACC DECLARE.
+
+module const
+  integer, parameter :: n = 100
+end module const
+
+subroutine check
+  use const
+
+  implicit none
+  integer i, x(n), y
+  common /BLOCK/ x, y
+
+  do i = 1, n
+     if (x(i) .ne. y) call abort
+  end do
+end subroutine check
+
+module m
+  use const
+  integer a(n), b
+  common /BLOCK/ a, b
+
+contains
+  subroutine mod_implicit_incr
+    implicit none
+    integer i
+
+    !$acc parallel loop
+    do i = 1, n
+       a(i) = b
+    end do
+    !$acc end parallel loop
+
+    call check
+  end subroutine mod_implicit_incr
+
+  subroutine mod_explicit_incr
+    implicit none
+    integer i
+
+    !$acc parallel loop copy(a(1:n)) copyin(b)
+    do i = 1, n
+       a(i) = b
+    end do
+    !$acc end parallel loop
+
+    call check
+  end subroutine mod_explicit_incr
+end module m
+
+subroutine sub_implicit_incr
+  use const
+
+  implicit none
+  integer i, x(n), y
+  common /BLOCK/ x, y
+
+  !$acc parallel loop
+  do i = 1, n
+     x(i) = y
+  end do
+  !$acc end parallel loop
+
+  call check
+end subroutine sub_implicit_incr
+
+subroutine sub_explicit_incr
+  use const
+
+  implicit none
+  integer i, x(n), y
+  common /BLOCK/ x, y
+
+  !$acc parallel loop copy(x(1:n)) copyin(y)
+  do i = 1, n
+     x(i) = y
+  end do
+  !$acc end parallel loop
+
+  call check
+end subroutine sub_explicit_incr
+
+program main
+  use m
+
+  implicit none
+
+  a(:) = -1
+  b = 5
+  call mod_implicit_incr
+
+  a(:) = -2
+  b = 6
+  call mod_explicit_incr
+
+  a(:) = -3
+  b = 7
+  call sub_implicit_incr
+
+  a(:) = -4
+  b = 8
+  call sub_explicit_incr
+end program main
Index: libgomp/testsuite/libgomp.oacc-fortran/common-block-2.f90
===================================================================
--- libgomp/testsuite/libgomp.oacc-fortran/common-block-2.f90	(nonexistent)
+++ libgomp/testsuite/libgomp.oacc-fortran/common-block-2.f90	(working copy)
@@ -0,0 +1,150 @@
+! Test data located inside common blocks.  This test does not execrise
+! ACC DECLARE.  All data clauses are explicit.
+
+module consts
+  integer, parameter :: n = 100
+end module consts
+
+subroutine validate
+  use consts
+
+  implicit none
+  integer i, j
+  real*4 x(n), y(n), z
+  common /BLOCK/ x, y, z, j
+
+  do i = 1, n
+     if (abs(x(i) - i - z) .ge. 0.0001) call abort
+  end do
+end subroutine validate
+
+subroutine incr
+  use consts
+
+  implicit none
+  integer i, j
+  real*4 x(n), y(n), z
+  common /BLOCK/ x, y, z, j
+
+  !$acc parallel loop pcopy(/BLOCK/)
+  do i = 1, n
+     x(i) = x(i) + z
+  end do
+  !$acc end parallel loop
+end subroutine incr
+
+program main
+  use consts
+
+  implicit none
+  integer i, j
+  real*4 a(n), b(n), c
+  common /BLOCK/ a, b, c, j
+
+  ! Test copyout, pcopy, device
+
+  !$acc data copyout(a, c)
+
+  c = 1.0
+
+  !$acc update device(c)
+
+  !$acc parallel loop pcopy(a)
+  do i = 1, n
+     a(i) = i
+  end do
+  !$acc end parallel loop
+
+  call incr
+  call incr
+  call incr
+  !$acc end data
+
+  c = 3.0
+  call validate
+
+  ! Test pcopy without copyout
+
+  c = 2.0
+  call incr
+  c = 5.0
+  call validate
+
+  ! Test create, delete, host, copyout, copyin
+
+  !$acc enter data create(b)
+
+  !$acc parallel loop pcopy(b)
+  do i = 1, n
+     b(i) = i
+  end do
+  !$acc end parallel loop
+
+  !$acc update host (b)
+
+  !$acc parallel loop pcopy(b) copyout(a) copyin(c)
+  do i = 1, n
+     a(i) = b(i) + c
+  end do
+  !$acc end parallel loop
+
+  !$acc exit data delete(b)
+
+  call validate
+
+  a(:) = b(:)
+  c = 0.0
+  call validate
+
+  ! Test copy
+
+  c = 1.0
+  !$acc parallel loop copy(/BLOCK/)
+  do i = 1, n
+     a(i) = b(i) + c
+  end do
+  !$acc end parallel loop
+
+  call validate
+
+  ! Test pcopyin, pcopyout FIXME
+
+  c = 2.0
+  !$acc data copyin(b, c) copyout(a)
+
+  !$acc parallel loop pcopyin(b, c) pcopyout(a)
+  do i = 1, n
+     a(i) = b(i) + c
+  end do
+  !$acc end parallel loop
+
+  !$acc end data
+
+  call validate
+
+  ! Test reduction, private
+
+  j = 0
+
+  !$acc parallel private(i) copy(j)
+  !$acc loop reduction(+:j)
+  do i = 1, n
+     j = j + 1
+  end do
+  !$acc end parallel
+
+  if (j .ne. n) call abort
+
+  ! Test firstprivate, copy
+
+  a(:) = 0
+  c = j
+
+  !$acc parallel loop firstprivate(c) copyout(a)
+  do i = 1, n
+     a(i) = i + c
+  end do
+  !$acc end parallel loop
+
+  call validate
+end program main
Index: libgomp/testsuite/libgomp.oacc-fortran/common-block-3.f90
===================================================================
--- libgomp/testsuite/libgomp.oacc-fortran/common-block-3.f90	(nonexistent)
+++ libgomp/testsuite/libgomp.oacc-fortran/common-block-3.f90	(working copy)
@@ -0,0 +1,137 @@
+! Test data located inside common blocks.  This test does not execrise
+! ACC DECLARE.  Most of the data clauses are implicit.
+
+module consts
+  integer, parameter :: n = 100
+end module consts
+
+subroutine validate
+  use consts
+
+  implicit none
+  integer i, j
+  real*4 x(n), y(n), z
+  common /BLOCK/ x, y, z, j
+
+  do i = 1, n
+     if (abs(x(i) - i - z) .ge. 0.0001) call abort
+  end do
+end subroutine validate
+
+subroutine incr_parallel
+  use consts
+
+  implicit none
+  integer i, j
+  real*4 x(n), y(n), z
+  common /BLOCK/ x, y, z, j
+
+  !$acc parallel loop
+  do i = 1, n
+     x(i) = x(i) + z
+  end do
+  !$acc end parallel loop
+end subroutine incr_parallel
+
+subroutine incr_kernels
+  use consts
+
+  implicit none
+  integer i, j
+  real*4 x(n), y(n), z
+  common /BLOCK/ x, y, z, j
+
+  !$acc kernels
+  do i = 1, n
+     x(i) = x(i) + z
+  end do
+  !$acc end kernels
+end subroutine incr_kernels
+
+program main
+  use consts
+
+  implicit none
+  integer i, j
+  real*4 a(n), b(n), c
+  common /BLOCK/ a, b, c, j
+
+  !$acc data copyout(a, c)
+
+  c = 1.0
+
+  !$acc update device(c)
+
+  !$acc parallel loop
+  do i = 1, n
+     a(i) = i
+  end do
+  !$acc end parallel loop
+
+  call incr_parallel
+  call incr_parallel
+  call incr_parallel
+  !$acc end data
+
+  c = 3.0
+  call validate
+
+  ! Test pcopy without copyout
+
+  c = 2.0
+  call incr_kernels
+  c = 5.0
+  call validate
+
+  !$acc kernels
+  do i = 1, n
+     b(i) = i
+  end do
+  !$acc end kernels
+
+  !$acc parallel loop
+  do i = 1, n
+     a(i) = b(i) + c
+  end do
+  !$acc end parallel loop
+
+  call validate
+
+  a(:) = b(:)
+  c = 0.0
+  call validate
+
+  ! Test copy
+
+  c = 1.0
+  !$acc parallel loop
+  do i = 1, n
+     a(i) = b(i) + c
+  end do
+  !$acc end parallel loop
+
+  call validate
+
+  c = 2.0
+  !$acc data copyin(b, c) copyout(a)
+
+  !$acc kernels
+  do i = 1, n
+     a(i) = b(i) + c
+  end do
+  !$acc end kernels
+
+  !$acc end data
+
+  call validate
+
+  j = 0
+
+  !$acc parallel loop reduction(+:j)
+  do i = 1, n
+     j = j + 1
+  end do
+  !$acc end parallel loop
+
+  if (j .ne. n) call abort
+end program main
Index: libgomp/testsuite/libgomp.oacc-fortran/deviceptr-1.f90
===================================================================
--- libgomp/testsuite/libgomp.oacc-fortran/deviceptr-1.f90	(nonexistent)
+++ libgomp/testsuite/libgomp.oacc-fortran/deviceptr-1.f90	(working copy)
@@ -0,0 +1,197 @@
+! { dg-do run }
+
+! Test the deviceptr clause with various directives
+! and in combination with other directives where
+! the deviceptr variable is implied.
+
+subroutine subr1 (a, b)
+  implicit none
+  integer, parameter :: N = 8
+  integer :: a(N)
+  integer :: b(N)
+  integer :: i = 0
+
+  !$acc data deviceptr (a)
+
+  !$acc parallel copy (b)
+    do i = 1, N
+      a(i) = i * 2
+      b(i) = a(i)
+    end do
+  !$acc end parallel
+
+  !$acc end data
+
+end subroutine
+
+subroutine subr2 (a, b)
+  implicit none
+  integer, parameter :: N = 8
+  integer :: a(N)
+  !$acc declare deviceptr (a)
+  integer :: b(N)
+  integer :: i = 0
+
+  !$acc parallel copy (b)
+    do i = 1, N
+      a(i) = i * 4
+      b(i) = a(i)
+    end do
+  !$acc end parallel
+
+end subroutine
+
+subroutine subr3 (a, b)
+  implicit none
+  integer, parameter :: N = 8
+  integer :: a(N)
+  !$acc declare deviceptr (a)
+  integer :: b(N)
+  integer :: i = 0
+
+  !$acc kernels copy (b)
+    do i = 1, N
+      a(i) = i * 8
+      b(i) = a(i)
+    end do
+  !$acc end kernels
+
+end subroutine
+
+subroutine subr4 (a, b)
+  implicit none
+  integer, parameter :: N = 8
+  integer :: a(N)
+  integer :: b(N)
+  integer :: i = 0
+
+  !$acc parallel deviceptr (a) copy (b)
+    do i = 1, N
+      a(i) = i * 16
+      b(i) = a(i)
+    end do
+  !$acc end parallel
+
+end subroutine
+
+subroutine subr5 (a, b)
+  implicit none
+  integer, parameter :: N = 8
+  integer :: a(N)
+  integer :: b(N)
+  integer :: i = 0
+
+  !$acc kernels deviceptr (a) copy (b)
+    do i = 1, N
+      a(i) = i * 32
+      b(i) = a(i)
+    end do
+  !$acc end kernels
+
+end subroutine
+
+subroutine subr6 (a, b)
+  implicit none
+  integer, parameter :: N = 8
+  integer :: a(N)
+  integer :: b(N)
+  integer :: i = 0
+
+  !$acc parallel deviceptr (a) copy (b)
+    do i = 1, N
+      b(i) = i
+    end do
+  !$acc end parallel
+
+end subroutine
+
+subroutine subr7 (a, b)
+  implicit none
+  integer, parameter :: N = 8
+  integer :: a(N)
+  integer :: b(N)
+  integer :: i = 0
+
+  !$acc data deviceptr (a)
+
+  !$acc parallel copy (b)
+    do i = 1, N
+      a(i) = i * 2
+      b(i) = a(i)
+    end do
+  !$acc end parallel
+
+  !$acc parallel copy (b)
+    do i = 1, N
+      a(i) = b(i) * 2
+      b(i) = a(i)
+    end do
+  !$acc end parallel
+
+  !$acc end data
+
+end subroutine
+
+program main
+  use iso_c_binding, only: c_ptr, c_f_pointer
+  implicit none
+  type (c_ptr) :: cp
+  integer, parameter :: N = 8
+  integer, pointer :: fp(:)
+  integer :: i = 0
+  integer :: b(N)
+
+  interface
+    function acc_malloc (s) bind (C)
+      use iso_c_binding, only: c_ptr, c_size_t
+      integer (c_size_t), value :: s
+      type (c_ptr) :: acc_malloc
+    end function
+  end interface
+
+  cp = acc_malloc (N * sizeof (fp(N)))
+  call c_f_pointer (cp, fp, [N])
+
+  call subr1 (fp, b)
+
+  do i = 1, N
+    if (b(i) .ne. i * 2) call abort
+  end do
+
+  call subr2 (fp, b)
+
+  do i = 1, N
+    if (b(i) .ne. i * 4) call abort
+  end do
+
+  call subr3 (fp, b)
+
+  do i = 1, N
+    if (b(i) .ne. i * 8) call abort
+  end do
+
+  call subr4 (fp, b)
+
+  do i = 1, N
+    if (b(i) .ne. i * 16) call abort
+  end do
+
+  call subr5 (fp, b)
+
+  do i = 1, N
+    if (b(i) .ne. i * 32) call abort
+  end do
+
+  call subr6 (fp, b)
+
+  do i = 1, N
+    if (b(i) .ne. i) call abort
+  end do
+
+  call subr7 (fp, b)
+
+  do i = 1, N
+    if (b(i) .ne. i * 4) call abort
+  end do
+
+end program main