diff mbox series

[OpenACC] Fix implicit mapping in enclosed 'acc data' constructions

Message ID d8061164-5386-1fd8-4e34-7104048d9426@codesourcery.com
State New
Headers show
Series [OpenACC] Fix implicit mapping in enclosed 'acc data' constructions | expand

Commit Message

Tobias Burnus Dec. 16, 2019, 5:31 p.m. UTC 
Currently, GCC uses *implicit 'copy'* for variables in a 
parallel/kernels construct, even though they are in a data clause of the 
enclosing *data construct*.

Example – will currently yield an an implicit 'copy(var)' in the 'parallel construct':
    … acc data copyout(var)
    … acc parallel
    var = 5

In the spec for OpenACC 2.6 or 2.7, one can find in the 'kernels 
Construction' and 'parallel Construction' towards the end of the 
'Description', the following wording:

"If there is no default(none) clause on the construct, the compiler will 
implicitly determine data attributes for variables that are referenced 
in the compute construct that do not have predetermined data attributes 
and do not appear in a data clause on the compute construct, a lexically 
containing data construct, or a visible declare directive. […] A scalar 
variable referenced in the kernels construct that does not appear in a 
data clause for the construct or any enclosing data construct will be 
treated as if it appeared in a copy clause."

(OpenACC 2+6 and 2.7 for kernels; for parallel, read as 
s/kernels/parallel/ and s/copy clause/firstprivate clause/ in the last 
sentence.)

That GCC uses a 'copy (or 'map(tofrom:') is an effect of the commit 
commit r230169 (GIT: b656be3a5e5d499ed45dba75943861a14d3aec55; Wed Nov 
11 14:24:09 2015 +0000), see 
https://gcc.gnu.org/ml/gcc-patches/2015-11/msg00809.html

The issue with the copy was noted by Thomas in 
https://gcc.gnu.org/ml/gcc-patches/2019-12/msg00147.html

This patch simply avoids the mapping in this case – while keeping the 
host data part.

OK for the trunk? Or have I missed some fine print/corner cases?

Cheers,

Tobias
diff mbox series

Patch

	gcc/
	* gimplify.c (omp_notice_variable): For OpenACC, don't imply copy clause
	for variables which appear in a data clause of the enclosing target/data
	context.
	
	gcc/testsuite/
	* c-c++-common/goacc/default-4.c: Update expected tree dump.
	* c-c++-common/goacc/deviceptr-4.c: Likewise.
	* gfortran.dg/goacc/default-4.f: Likewise.

	libgomp/
	* testsuite/libgomp.oacc-fortran/implicit-mapping.F90: New.

diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index 9073680cb31..22e92393e7f 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -7471,37 +7471,48 @@  omp_notice_variable (struct gimplify_omp_ctx *ctx, tree decl, bool in_code)
 	  if ((ctx->region_type & ORT_ACC) && octx)
 	    {
 	      /* Look in outer OpenACC contexts, to see if there's a
 		 data attribute for this variable.  */
 	      omp_notice_variable (octx, decl, in_code);
 
 	      for (; octx; octx = octx->outer_context)
 		{
 		  if (!(octx->region_type & (ORT_TARGET_DATA | ORT_TARGET)))
 		    break;
 		  splay_tree_node n2
 		    = splay_tree_lookup (octx->variables,
 					 (splay_tree_key) decl);
 		  if (n2)
 		    {
 		      if (octx->region_type == ORT_ACC_HOST_DATA)
 		        error ("variable %qE declared in enclosing "
 			       "%<host_data%> region", DECL_NAME (decl));
-		      nflags |= GOVD_MAP;
+		      /* Already mapped in outer context; hence, no implicit
+			 mapping. For 'acc data', we still need to add a
+			 present() to associate the identifier with the
+			 device data. */
+		      if (octx->region_type & ORT_TARGET
+			  && n2->value & (GOVD_FIRSTPRIVATE | GOVD_MAP))
+			return ret;
+		      if (octx->region_type & ORT_TARGET_DATA
+			  && n2->value & GOVD_MAP)
+			nflags |= GOVD_MAP | GOVD_MAP_FORCE_PRESENT;
+		      else
+			nflags |= GOVD_MAP;
 		      if (octx->region_type == ORT_ACC_DATA
 			  && (n2->value & GOVD_MAP_0LEN_ARRAY))
 			nflags |= GOVD_MAP_0LEN_ARRAY;
 		      goto found_outer;
 		    }
 		}
 	    }
 
 	  if ((nflags & ~(GOVD_MAP_TO_ONLY | GOVD_MAP_FROM_ONLY
 			  | GOVD_MAP_ALLOC_ONLY)) == flags)
 	    {
 	      tree type = TREE_TYPE (decl);
 
 	      if (gimplify_omp_ctxp->target_firstprivatize_array_bases
 		  && lang_hooks.decls.omp_privatize_by_reference (decl))
 		type = TREE_TYPE (type);
 	      if (!lang_hooks.types.omp_mappable_type (type))
 		{
diff --git a/gcc/testsuite/c-c++-common/goacc/default-4.c b/gcc/testsuite/c-c++-common/goacc/default-4.c
index 867175d4847..76deaf36d25 100644
--- a/gcc/testsuite/c-c++-common/goacc/default-4.c
+++ b/gcc/testsuite/c-c++-common/goacc/default-4.c
@@ -1,66 +1,66 @@ 
 /* OpenACC default clause inside data construct.  */
 
 /* { dg-additional-options "-fdump-tree-gimple" } */
 
 void f1 ()
 {
   int f1_a = 2;
   float f1_b[2];
 
 #pragma acc data copyin (f1_a) copyout (f1_b)
   /* { dg-final { scan-tree-dump-times "omp target oacc_data map\\(from:f1_b \[^\\)\]+\\) map\\(to:f1_a" 1 "gimple" } } */
   {
 #pragma acc kernels
-    /* { dg-final { scan-tree-dump-times "omp target oacc_kernels map\\(tofrom:f1_b \[^\\)\]+\\) map\\(tofrom:f1_a" 1 "gimple" } } */
+    /* { dg-final { scan-tree-dump-times "omp target oacc_kernels map\\(force_present:f1_b \[^\\)\]+\\) map\\(force_present:f1_a" 1 "gimple" } } */
     {
       f1_b[0] = f1_a;
     }
 #pragma acc parallel
-    /* { dg-final { scan-tree-dump-times "omp target oacc_parallel map\\(tofrom:f1_b \[^\\)\]+\\) map\\(tofrom:f1_a" 1 "gimple" } } */
+    /* { dg-final { scan-tree-dump-times "omp target oacc_parallel map\\(force_present:f1_b \[^\\)\]+\\) map\\(force_present:f1_a" 1 "gimple" } } */
     {
       f1_b[0] = f1_a;
     }
   }
 }
 
 void f2 ()
 {
   int f2_a = 2;
   float f2_b[2];
 
 #pragma acc data copyin (f2_a) copyout (f2_b)
   /* { dg-final { scan-tree-dump-times "omp target oacc_data map\\(from:f2_b \[^\\)\]+\\) map\\(to:f2_a" 1 "gimple" } } */
   {
 #pragma acc kernels default (none)
-    /* { dg-final { scan-tree-dump-times "omp target oacc_kernels default\\(none\\) map\\(tofrom:f2_b \[^\\)\]+\\) map\\(tofrom:f2_a" 1 "gimple" } } */
+    /* { dg-final { scan-tree-dump-times "omp target oacc_kernels default\\(none\\) map\\(force_present:f2_b \[^\\)\]+\\) map\\(force_present:f2_a" 1 "gimple" } } */
     {
       f2_b[0] = f2_a;
     }
 #pragma acc parallel default (none)
-    /* { dg-final { scan-tree-dump-times "omp target oacc_parallel default\\(none\\) map\\(tofrom:f2_b \[^\\)\]+\\) map\\(tofrom:f2_a" 1 "gimple" } } */
+    /* { dg-final { scan-tree-dump-times "omp target oacc_parallel default\\(none\\) map\\(force_present:f2_b \[^\\)\]+\\) map\\(force_present:f2_a" 1 "gimple" } } */
     {
       f2_b[0] = f2_a;
     }
   }
 }
 
 void f3 ()
 {
   int f3_a = 2;
   float f3_b[2];
 
 #pragma acc data copyin (f3_a) copyout (f3_b)
   /* { dg-final { scan-tree-dump-times "omp target oacc_data map\\(from:f3_b \[^\\)\]+\\) map\\(to:f3_a" 1 "gimple" } } */
   {
 #pragma acc kernels default (present)
-    /* { dg-final { scan-tree-dump-times "omp target oacc_kernels default\\(present\\) map\\(tofrom:f3_b \[^\\)\]+\\) map\\(tofrom:f3_a" 1 "gimple" } } */
+    /* { dg-final { scan-tree-dump-times "omp target oacc_kernels default\\(present\\) map\\(force_present:f3_b \[^\\)\]+\\) map\\(force_present:f3_a" 1 "gimple" } } */
     {
       f3_b[0] = f3_a;
     }
 #pragma acc parallel default (present)
-    /* { dg-final { scan-tree-dump-times "omp target oacc_parallel default\\(present\\) map\\(tofrom:f3_b \[^\\)\]+\\) map\\(tofrom:f3_a" 1 "gimple" } } */
+    /* { dg-final { scan-tree-dump-times "omp target oacc_parallel default\\(present\\) map\\(force_present:f3_b \[^\\)\]+\\) map\\(force_present:f3_a" 1 "gimple" } } */
     {
       f3_b[0] = f3_a;
     }
   }
 }
diff --git a/gcc/testsuite/c-c++-common/goacc/deviceptr-4.c b/gcc/testsuite/c-c++-common/goacc/deviceptr-4.c
index db1b91633a6..7962843ba49 100644
--- a/gcc/testsuite/c-c++-common/goacc/deviceptr-4.c
+++ b/gcc/testsuite/c-c++-common/goacc/deviceptr-4.c
@@ -1,11 +1,11 @@ 
 /* { dg-additional-options "-fdump-tree-gimple" } */
 
 void
 subr (int *a)
 {
 #pragma acc data deviceptr (a)
 #pragma acc parallel
   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_present:a" 1 "gimple" } } */
diff --git a/gcc/testsuite/gfortran.dg/goacc/default-4.f b/gcc/testsuite/gfortran.dg/goacc/default-4.f
index 30f411f70ab..ce8fde422db 100644
--- a/gcc/testsuite/gfortran.dg/goacc/default-4.f
+++ b/gcc/testsuite/gfortran.dg/goacc/default-4.f
@@ -1,57 +1,57 @@ 
 ! OpenACC default clause inside data construct.
 
 ! { dg-additional-options "-fdump-tree-gimple" } 
 
       SUBROUTINE F1
       IMPLICIT NONE
       INTEGER :: F1_A = 2
       REAL, DIMENSION (2) :: F1_B
 
 !$ACC DATA COPYIN (F1_A) COPYOUT (F1_B)
 ! { dg-final { scan-tree-dump-times "omp target oacc_data map\\(to:f1_a \[^\\)\]+\\) map\\(from:f1_b" 1 "gimple" } }
 !$ACC KERNELS
-! { dg-final { scan-tree-dump-times "omp target oacc_kernels map\\(tofrom:f1_b \[^\\)\]+\\) map\\(tofrom:f1_a" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "omp target oacc_kernels map\\(force_present:f1_b \[^\\)\]+\\) map\\(force_present:f1_a" 1 "gimple" } }
       F1_B(1) = F1_A;
 !$ACC END KERNELS
 !$ACC PARALLEL
-! { dg-final { scan-tree-dump-times "omp target oacc_parallel map\\(tofrom:f1_b \[^\\)\]+\\) map\\(tofrom:f1_a" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "omp target oacc_parallel map\\(force_present:f1_b \[^\\)\]+\\) map\\(force_present:f1_a" 1 "gimple" } }
       F1_B(1) = F1_A;
 !$ACC END PARALLEL
 !$ACC END DATA
       END SUBROUTINE F1
 
       SUBROUTINE F2
       IMPLICIT NONE
       INTEGER :: F2_A = 2
       REAL, DIMENSION (2) :: F2_B
 
 !$ACC DATA COPYIN (F2_A) COPYOUT (F2_B)
 ! { dg-final { scan-tree-dump-times "omp target oacc_data map\\(to:f2_a \[^\\)\]+\\) map\\(from:f2_b" 1 "gimple" } }
 !$ACC KERNELS DEFAULT (NONE)
-! { dg-final { scan-tree-dump-times "omp target oacc_kernels default\\(none\\) map\\(tofrom:f2_b \[^\\)\]+\\) map\\(tofrom:f2_a" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "omp target oacc_kernels default\\(none\\) map\\(force_present:f2_b \[^\\)\]+\\) map\\(force_present:f2_a" 1 "gimple" } }
       F2_B(1) = F2_A;
 !$ACC END KERNELS
 !$ACC PARALLEL DEFAULT (NONE)
-! { dg-final { scan-tree-dump-times "omp target oacc_parallel default\\(none\\) map\\(tofrom:f2_b \[^\\)\]+\\) map\\(tofrom:f2_a" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "omp target oacc_parallel default\\(none\\) map\\(force_present:f2_b \[^\\)\]+\\) map\\(force_present:f2_a" 1 "gimple" } }
       F2_B(1) = F2_A;
 !$ACC END PARALLEL
 !$ACC END DATA
       END SUBROUTINE F2
 
       SUBROUTINE F3
       IMPLICIT NONE
       INTEGER :: F3_A = 2
       REAL, DIMENSION (2) :: F3_B
 
 !$ACC DATA COPYIN (F3_A) COPYOUT (F3_B)
 ! { dg-final { scan-tree-dump-times "omp target oacc_data map\\(to:f3_a \[^\\)\]+\\) map\\(from:f3_b" 1 "gimple" } }
 !$ACC KERNELS DEFAULT (PRESENT)
-! { dg-final { scan-tree-dump-times "omp target oacc_kernels default\\(present\\) map\\(tofrom:f3_b \[^\\)\]+\\) map\\(tofrom:f3_a" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "omp target oacc_kernels default\\(present\\) map\\(force_present:f3_b \[^\\)\]+\\) map\\(force_present:f3_a" 1 "gimple" } }
       F3_B(1) = F3_A;
 !$ACC END KERNELS
 !$ACC PARALLEL DEFAULT (PRESENT)
-! { dg-final { scan-tree-dump-times "omp target oacc_parallel default\\(present\\) map\\(tofrom:f3_b \[^\\)\]+\\) map\\(tofrom:f3_a" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "omp target oacc_parallel default\\(present\\) map\\(force_present:f3_b \[^\\)\]+\\) map\\(force_present:f3_a" 1 "gimple" } }
       F3_B(1) = F3_A;
 !$ACC END PARALLEL
 !$ACC END DATA
       END SUBROUTINE F3
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/implicit-mapping.F90 b/libgomp/testsuite/libgomp.oacc-fortran/implicit-mapping.F90
new file mode 100644
index 00000000000..2e45d5014b9
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/implicit-mapping.F90
@@ -0,0 +1,179 @@ 
+! { dg-do run }
+! { dg-additional-options "-fdump-tree-omplower" }
+!
+! Check that implicit mapping happens as follows:
+! - parallel: firstprivate for scalars and
+!             for arrays/aggregated types: copy - or with default(present) - present
+! - kernels:  copy for scalars and
+!             for arrays/aggregated types: copy - or with default(present) - present
+! - No implicit mapping if 'predetermined' or 'do not appear in a data clause
+!   on the compute construct, a lexically containing data construct'
+!
+! Exact wording, see end of 'Description' in the 'kernels Construct' or
+! the 'parallel Construct' section of OpenACC 2.6 or 2.7.
+!
+implicit none
+integer :: var1, var2, var3, var4, arr1(3), arr2(3), arr3(3), arr4(3)
+integer :: other1, other2, other3, other4, vec1(3), vec2(3), vec3(3), vec4(3)
+integer :: varout2, otherout2
+
+var1 = 122
+var3 = 135
+arr1 = [14,325,35]
+arr3 = [24,221,73]
+
+var2 = 0
+var4 = 0
+varout2 = 0
+
+!$acc data copyin(var1, var3, arr1, arr3) copyout(varout2)
+#if !ACC_MEM_SHARED
+var1 = 3534
+var3 = 26
+arr1 = [13,54,216]
+arr3 = [4,5231,5253]
+#endif
+
+!$acc parallel   ! Expect: 'firstprivate(var2)' copy(arr2) and no mapping of 'var1' or 'arr1'
+var2 = var1
+arr2 = arr1
+varout2 = var2
+!$acc end parallel
+
+!$acc kernels   ! Expect: 'copy(var4, arr4)' and no mapping of 'var3' or 'arr3'
+var4 = var3
+arr4 = arr3
+!$acc end kernels
+!$acc end data
+
+if (varout2 /= 122) stop 1
+if (any (arr2 /= [14,325,35])) stop 2
+if (var4 /= 135) stop 3
+if (any (arr4 /= [24,221,73])) stop 4
+if (var2 /= 0) stop 5 ! due to first private
+
+#if !ACC_MEM_SHARED
+if (var1 /= 3534) stop 6
+if (var3 /= 26) stop 7
+if (any (arr1 /= [13,54,216])) stop 8
+if (any (arr3 /= [4,5231,5253])) stop 9
+#else
+if (var1 /= varout2) stop 10
+if (var1 /= 122) stop 11
+if (var3 /= var4) stop 12
+if (any (arr1 /= arr2)) stop 13
+if (any (arr3 /= arr4)) stop 14
+#endif
+
+! -------------------------
+
+other1 = 25
+other3 = 573
+vec1 = [4,56,17]
+vec3 = [62,87,593]
+other2 = 0
+other4 = 0
+vec2 = 0
+vec4 = 0
+otherout2 = 0
+
+!$acc enter data create(vec2, vec4)
+
+!$acc data copyin(other1, other3, vec1, vec3) copyout(otherout2)
+#if !ACC_MEM_SHARED
+other1 = 465
+other3 = 6
+vec1 = [754,412,562]
+vec3 = [58,48,671]
+#endif
+
+!$acc parallel default(present)   ! Expect: 'firstprivate(other2) present(vec2)', only
+other2 = other1
+vec2 = vec1
+otherout2 = other2
+!$acc end parallel
+
+!$acc kernels default(present)   ! Expect: 'copy(other4) present(vec4)', only
+other4 = other3
+vec4 = vec3
+!$acc end kernels
+!$acc end data
+
+#if !ACC_MEM_SHARED
+if (any (vec2 /= 0)) stop 21
+if (any (vec4 /= 0)) stop 22
+#endif
+if (otherout2 /= 25) stop 23
+if (other4 /= 573) stop 24
+if (other2 /= 0) stop 25 ! due to first private
+
+!$acc exit data copyout(vec2, vec4)
+if (any (vec2 /= [4,56,17])) stop 26
+if (any (vec4 /= [62,87,593])) stop 27
+
+#if !ACC_MEM_SHARED
+if (other1 /= 465) stop 28
+if (other3 /= 6) stop 29
+if (any (vec1 /= [754,412,562])) stop 30
+if (any (vec3 /= [58,48,671])) stop 31
+#else
+if (other1 /= otherout2) stop 32
+if (other1 /= 25) stop 33
+if (other3 /= other4) stop 34
+if (any (vec1 /= vec2)) stop 35
+if (any (vec3 /= vec4)) stop 36
+#endif
+
+end
+
+! Dump for  acc data copyin(var1, var3, arr1, arr3) copyout(varout2)
+! { dg-final { scan-tree-dump-times "oacc_data \[^\\n\]*map\\(to:var1 \\\[len: 4\\\]\\)" 1 "omplower" } }
+! { dg-final { scan-tree-dump-times "oacc_data \[^\\n\]*map\\(to:arr1 \\\[len: 12\\\]\\)" 1 "omplower" } }
+! { dg-final { scan-tree-dump-times "oacc_data \[^\\n\]*map\\(from:varout2 \\\[len: 4\\\]\\)" 1 "omplower" } }
+! { dg-final { scan-tree-dump-not "oacc_data \[^\\n\]*var2" "omplower" } }
+! { dg-final { scan-tree-dump-not "oacc_data \[^\\n\]*var4" "omplower" } }
+! { dg-final { scan-tree-dump-not "oacc_data \[^\\n\]*arr2" "omplower" } }
+! { dg-final { scan-tree-dump-not "oacc_data \[^\\n\]*arr4" "omplower" } }
+
+! Dump for first  acc parallel
+! { dg-final { scan-tree-dump-times "oacc_parallel \[^\\n\]*firstprivate\\(var2\\)" 1 "omplower" } }
+! { dg-final { scan-tree-dump-times "oacc_parallel \[^\\n\]*map\\(tofrom:arr2 \\\[len: 12\\\]\\)" 1 "omplower" } }
+! { dg-final { scan-tree-dump-times "oacc_parallel \[^\\n\]*map\\(force_present:var1 \\\[len: 4\\\]\\)" 1 "omplower" } }
+! { dg-final { scan-tree-dump-times "oacc_parallel \[^\\n\]*map\\(force_present:arr1 \\\[len: 12\\\]\\)" 1 "omplower" } }
+! { dg-final { scan-tree-dump-not "oacc_parallel \[^\\n\]*var3" "omplower" } }
+! { dg-final { scan-tree-dump-not "oacc_parallel \[^\\n\]*arr3" "omplower" } }
+
+! Dump for first  acc kernels
+! { dg-final { scan-tree-dump-times "oacc_kernels \[^\\n\]*map\\(force_tofrom:var4 \\\[len: 4\\\]\\)" 1 "omplower" } }
+! { dg-final { scan-tree-dump-times "oacc_kernels \[^\\n\]*map\\(tofrom:arr4 \\\[len: 12\\\]\\)" 1 "omplower" } }
+! { dg-final { scan-tree-dump-times "oacc_kernels \[^\\n\]*map\\(force_present:var3 \\\[len: 4\\\]\\)" 1 "omplower" } }
+! { dg-final { scan-tree-dump-times "oacc_kernels \[^\\n\]*map\\(force_present:arr3 \\\[len: 12\\\]\\)" 1 "omplower" } }
+! { dg-final { scan-tree-dump-not "oacc_kernels \[^\\n\]*var1" "omplower" } }
+! { dg-final { scan-tree-dump-not "oacc_kernels \[^\\n\]*arr1" "omplower" } }
+
+! ---------------
+
+! Dump for  acc data copyin(other1, other3, vec1, vec3) copyout(otherout2)
+! { dg-final { scan-tree-dump-times "oacc_data \[^\\n\]*map\\(to:other1 \\\[len: 4\\\]\\)" 1 "omplower" } }
+! { dg-final { scan-tree-dump-times "oacc_data \[^\\n\]*map\\(to:vec1 \\\[len: 12\\\]\\)" 1 "omplower" } }
+! { dg-final { scan-tree-dump-times "oacc_data \[^\\n\]*map\\(from:otherout2 \\\[len: 4\\\]\\)" 1 "omplower" } }
+! { dg-final { scan-tree-dump-not "oacc_data \[^\\n\]*other2" "omplower" } }
+! { dg-final { scan-tree-dump-not "oacc_data \[^\\n\]*other4" "omplower" } }
+! { dg-final { scan-tree-dump-not "oacc_data \[^\\n\]*vec2" "omplower" } }
+! { dg-final { scan-tree-dump-not "oacc_data \[^\\n\]*vec4" "omplower" } }
+
+! Dump for first  acc parallel
+! { dg-final { scan-tree-dump-times "oacc_parallel \[^\\n\]*firstprivate\\(other2\\)" 1 "omplower" } }
+! { dg-final { scan-tree-dump-times "oacc_parallel \[^\\n\]*map\\(force_present:vec2 \\\[len: 12\\\]\\)" 1 "omplower" } }
+! { dg-final { scan-tree-dump-times "oacc_parallel \[^\\n\]*map\\(force_present:other1 \\\[len: 4\\\]\\)" 1 "omplower" } }
+! { dg-final { scan-tree-dump-times "oacc_parallel \[^\\n\]*map\\(force_present:vec1 \\\[len: 12\\\]\\)" 1 "omplower" } }
+! { dg-final { scan-tree-dump-not "oacc_parallel \[^\\n\]*other3" "omplower" } }
+! { dg-final { scan-tree-dump-not "oacc_parallel \[^\\n\]*vec3" "omplower" } }
+
+! Dump for first  acc kernels
+! { dg-final { scan-tree-dump-times "oacc_kernels \[^\\n\]*map\\(force_tofrom:other4 \\\[len: 4\\\]\\)" 1 "omplower" } }
+! { dg-final { scan-tree-dump-times "oacc_kernels \[^\\n\]*map\\(force_present:vec4 \\\[len: 12\\\]\\)" 1 "omplower" } }
+! { dg-final { scan-tree-dump-times "oacc_kernels \[^\\n\]*map\\(force_present:other3 \\\[len: 4\\\]\\)" 1 "omplower" } }
+! { dg-final { scan-tree-dump-times "oacc_kernels \[^\\n\]*map\\(force_present:vec3 \\\[len: 12\\\]\\)" 1 "omplower" } }
+! { dg-final { scan-tree-dump-not "oacc_kernels \[^\\n\]*other1" "omplower" } }
+! { dg-final { scan-tree-dump-not "oacc_kernels \[^\\n\]*vec1" "omplower" } }