fortran: OpenMP/OpenACC array mapping alignment fix (PR90030)

Fix issue with the Fortran front-end when mapping arrays: when creating the
data MEM_REF for the map clause, there was a convention of casting the
referencing pointer to 'c_char *' by
fold_convert (build_pointer_type (char_type_node), ptr).

This causes the alignment passed to the libgomp runtime for array data
hardwared to '1', and causes alignment errors on the offload target.

This patch fixes this by removing the char_type_node pointer converts, and
adding gcc_asserts to ensure POINTER_TYPE_P (TREE_TYPE (ptr)).

	PR fortran/90030

gcc/fortran/ChangeLog:

	* trans-openmp.c (gfc_omp_finish_clause): Remove fold_convert to pointer
	to char_type_node, add gcc_assert of POINTER_TYPE_P.
	(gfc_trans_omp_array_section): Likewise.
	(gfc_trans_omp_clauses): Likewise.

gcc/testsuite/ChangeLog:

	* gfortran.dg/goacc/finalize-1.f: Adjust scan test.
	* gfortran.dg/gomp/affinity-clause-1.f90: Likewise.
	* gfortran.dg/gomp/affinity-clause-5.f90: Likewise.
	* gfortran.dg/gomp/defaultmap-4.f90: Likewise.
	* gfortran.dg/gomp/defaultmap-5.f90: Likewise.
	* gfortran.dg/gomp/defaultmap-6.f90: Likewise.
	* gfortran.dg/gomp/map-3.f90: Likewise.
	* gfortran.dg/gomp/pr78260-2.f90: Likewise.
	* gfortran.dg/gomp/pr78260-3.f90: Likewise.

libgomp/ChangeLog:

	* testsuite/libgomp.oacc-fortran/pr90030.f90: New test.
	* testsuite/libgomp.fortran/pr90030.f90: New test.

diff --git a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c
index 5b3c310..18268fb 100644 (file)
--- a/gcc/fortran/trans-openmp.c
+++ b/gcc/fortran/trans-openmp.c
@@ -1564,7 +1564,7 @@ gfc_omp_finish_clause (tree c, gimple_seq *pre_p, bool openacc)
       if (present)
        ptr = gfc_build_cond_assign_expr (&block, present, ptr,
                                          null_pointer_node);
-      ptr = fold_convert (build_pointer_type (char_type_node), ptr);
+      gcc_assert (POINTER_TYPE_P (TREE_TYPE (ptr)));
       ptr = build_fold_indirect_ref (ptr);
       OMP_CLAUSE_DECL (c) = ptr;
       c2 = build_omp_clause (input_location, OMP_CLAUSE_MAP);
@@ -2381,7 +2381,7 @@ gfc_trans_omp_array_section (stmtblock_t *block, gfc_omp_namelist *n,
                                            OMP_CLAUSE_SIZE (node), elemsz);
     }
   gcc_assert (se.post.head == NULL_TREE);
-  ptr = fold_convert (build_pointer_type (char_type_node), ptr);
+  gcc_assert (POINTER_TYPE_P (TREE_TYPE (ptr)));
   OMP_CLAUSE_DECL (node) = build_fold_indirect_ref (ptr);
   ptr = fold_convert (ptrdiff_type_node, ptr);
 
@@ -2849,8 +2849,7 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
                  if (GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (decl)))
                    {
                      decl = gfc_conv_descriptor_data_get (decl);
-                     decl = fold_convert (build_pointer_type (char_type_node),
-                                          decl);
+                     gcc_assert (POINTER_TYPE_P (TREE_TYPE (decl)));
                      decl = build_fold_indirect_ref (decl);
                    }
                  else if (DECL_P (decl))
@@ -2873,8 +2872,7 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
                    }
                  gfc_add_block_to_block (&iter_block, &se.pre);
                  gfc_add_block_to_block (&iter_block, &se.post);
-                 ptr = fold_convert (build_pointer_type (char_type_node),
-                                     ptr);
+                 gcc_assert (POINTER_TYPE_P (TREE_TYPE (ptr)));
                  OMP_CLAUSE_DECL (node) = build_fold_indirect_ref (ptr);
                }
              if (list == OMP_LIST_DEPEND)
@@ -3117,8 +3115,7 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
                      if (present)
                        ptr = gfc_build_cond_assign_expr (block, present, ptr,
                                                          null_pointer_node);
-                     ptr = fold_convert (build_pointer_type (char_type_node),
-                                         ptr);
+                     gcc_assert (POINTER_TYPE_P (TREE_TYPE (ptr)));
                      ptr = build_fold_indirect_ref (ptr);
                      OMP_CLAUSE_DECL (node) = ptr;
                      node2 = build_omp_clause (input_location,
@@ -3555,8 +3552,7 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
                    {
                      tree type = TREE_TYPE (decl);
                      tree ptr = gfc_conv_descriptor_data_get (decl);
-                     ptr = fold_convert (build_pointer_type (char_type_node),
-                                         ptr);
+                     gcc_assert (POINTER_TYPE_P (TREE_TYPE (ptr)));
                      ptr = build_fold_indirect_ref (ptr);
                      OMP_CLAUSE_DECL (node) = ptr;
                      OMP_CLAUSE_SIZE (node)
@@ -3606,8 +3602,7 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
                                       OMP_CLAUSE_SIZE (node), elemsz);
                    }
                  gfc_add_block_to_block (block, &se.post);
-                 ptr = fold_convert (build_pointer_type (char_type_node),
-                                     ptr);
+                 gcc_assert (POINTER_TYPE_P (TREE_TYPE (ptr)));
                  OMP_CLAUSE_DECL (node) = build_fold_indirect_ref (ptr);
                }
              omp_clauses = gfc_trans_add_clause (node, omp_clauses);

diff --git a/gcc/testsuite/gfortran.dg/goacc/finalize-1.f b/gcc/testsuite/gfortran.dg/goacc/finalize-1.f
index b706b38..1e5bf0b 100644 (file)
--- a/gcc/testsuite/gfortran.dg/goacc/finalize-1.f
+++ b/gcc/testsuite/gfortran.dg/goacc/finalize-1.f
@@ -20,8 +20,8 @@
 ! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_exit_data map\\(delete:del_f \\\[len: \[0-9\]+\\\]\\) finalize$" 1 "gimple" } }
 
 !$ACC EXIT DATA FINALIZE DELETE (del_f_p(2:5))
-! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(release:\\*\\(c_char \\*\\) parm\\.0\\.data \\\[len: \[^\\\]\]+\\\]\\) map\\(to:del_f_p \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(alloc:\\(integer\\(kind=1\\)\\\[0:\\\] \\* restrict\\) del_f_p\\.data \\\[pointer assign, bias: \\(.*int.*\\) parm\\.0\\.data - \\(.*int.*\\) del_f_p\\.data\\\]\\) finalize;$" 1 "original" } }
-! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_exit_data map\\(delete:MEM\\\[\\(c_char \\*\\)\[^\\\]\]+\\\] \\\[len: \[^\\\]\]+\\\]\\) map\\(to:del_f_p \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(alloc:del_f_p\\.data \\\[pointer assign, bias: \[^\\\]\]+\\\]\\) finalize$" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(release:\\*\\(integer\\(kind=.\\)\\\[0:\\\] \\*\\) parm\\.0\\.data \\\[len: \[^\\\]\]+\\\]\\) map\\(to:del_f_p \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(alloc:\\(integer\\(kind=1\\)\\\[0:\\\] \\* restrict\\) del_f_p\\.data \\\[pointer assign, bias: \\(.*int.*\\) parm\\.0\\.data - \\(.*int.*\\) del_f_p\\.data\\\]\\) finalize;$" 1 "original" } }
+! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_exit_data map\\(delete:MEM <\[^>\]+> \\\[\\(integer\\(kind=.\\)\\\[0:\\\] \\*\\)_\[0-9\]+\\\] \\\[len: \[^\\\]\]+\\\]\\) map\\(to:del_f_p \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(alloc:del_f_p\\.data \\\[pointer assign, bias: \[^\\\]\]+\\\]\\) finalize$" 1 "gimple" } }
 
 !$ACC EXIT DATA COPYOUT (cpo_r)
 ! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(from:cpo_r\\);$" 1 "original" } }
@@ -32,6 +32,6 @@
 ! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_exit_data map\\(force_from:cpo_f \\\[len: \[0-9\]+\\\]\\) finalize$" 1 "gimple" } }
 
 !$ACC EXIT DATA COPYOUT (cpo_f_p(4:10)) FINALIZE
-! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(from:\\*\\(c_char \\*\\) parm\\.1\\.data \\\[len: \[^\\\]\]+\\\]\\) map\\(to:cpo_f_p \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(alloc:\\(integer\\(kind=1\\)\\\[0:\\\] \\* restrict\\) cpo_f_p\\.data \\\[pointer assign, bias: \\(.*int.*\\) parm\\.1\\.data - \\(.*int.*\\) cpo_f_p\\.data\\\]\\) finalize;$" 1 "original" } }
-! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_exit_data map\\(force_from:MEM\\\[\\(c_char \\*\\)\[^\\\]\]+\\\] \\\[len: \[^\\\]\]+\\\]\\) map\\(to:cpo_f_p \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(alloc:cpo_f_p\\.data \\\[pointer assign, bias: \[^\\\]\]+\\\]\\) finalize$" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(from:\\*\\(integer\\(kind=.\\)\\\[0:\\\] \\*\\) parm\\.1\\.data \\\[len: \[^\\\]\]+\\\]\\) map\\(to:cpo_f_p \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(alloc:\\(integer\\(kind=1\\)\\\[0:\\\] \\* restrict\\) cpo_f_p\\.data \\\[pointer assign, bias: \\(.*int.*\\) parm\\.1\\.data - \\(.*int.*\\) cpo_f_p\\.data\\\]\\) finalize;$" 1 "original" } }
+! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_exit_data map\\(force_from:MEM <\[^>\]+> \\\[\\(integer\\(kind=.\\)\\\[0:\\\] \\*\\)_\[0-9\]+\\\] \\\[len: \[^\\\]\]+\\\]\\) map\\(to:cpo_f_p \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(alloc:cpo_f_p\\.data \\\[pointer assign, bias: \[^\\\]\]+\\\]\\) finalize$" 1 "gimple" } }
       END SUBROUTINE f

diff --git a/gcc/testsuite/gfortran.dg/gomp/affinity-clause-1.f90 b/gcc/testsuite/gfortran.dg/gomp/affinity-clause-1.f90
index 13bdd36..08c7740 100644 (file)
--- a/gcc/testsuite/gfortran.dg/gomp/affinity-clause-1.f90
+++ b/gcc/testsuite/gfortran.dg/gomp/affinity-clause-1.f90
@@ -22,12 +22,12 @@ end
 
 ! { dg-final { scan-tree-dump-times "D\\.\[0-9\]+ = .integer.kind=4.. __builtin_cosf ..real.kind=4.. a \\+ 1.0e\\+0\\);" 2 "original" } }
 
-! { dg-final { scan-tree-dump-times "#pragma omp task affinity\\(iterator\\(integer\\(kind=4\\) jj=2:5:2, integer\\(kind=4\\) i=D\\.\[0-9\]+:5:1\\):\\*\\(c_char \\*\\) &b\\\[.* <?i>? \\+ -1\\\]\\) affinity\\(iterator\\(integer\\(kind=4\\) jj=2:5:2, integer\\(kind=4\\) i=D\\.\[0-9\]+:5:1\\):\\*\\(c_char \\*\\) &d\\\[\\(.*jj \\* 5 \\+ .* <?i>?\\) \\+ -6\\\]\\)" 1 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma omp task affinity\\(iterator\\(integer\\(kind=4\\) jj=2:5:2, integer\\(kind=4\\) i=D\\.\[0-9\]+:5:1\\):b\\\[.* <?i>? \\+ -1\\\]\\) affinity\\(iterator\\(integer\\(kind=4\\) jj=2:5:2, integer\\(kind=4\\) i=D\\.\[0-9\]+:5:1\\):d\\\[\\(.*jj \\* 5 \\+ .* <?i>?\\) \\+ -6\\\]\\)" 1 "original" } }
 
-! { dg final { scan-tree-dump-times "#pragma omp task affinity\\(iterator\\(integer\\(kind=4\\) i=D.3938:5:1\\):\\*\\(c_char \\*\\) &b\\\[\\(.* <?i>? \\+ -1\\\]\\) affinity\\(iterator\\(integer\\(kind=4\\) i=D\\.\[0-9\]+:5:1\\):\\*\\(c_char \\*\\) &d\\\[\\(\\(integer\\(kind=8\\)\\) i \\+ -1\\) \\* 6\\\]\\)"  1 "original" } }
+! { dg final { scan-tree-dump-times "#pragma omp task affinity\\(iterator\\(integer\\(kind=4\\) i=D\\.\[0-9\]+:5:1\\):b\\\[\\(.* <?i>? \\+ -1\\\]\\) affinity\\(iterator\\(integer\\(kind=4\\) i=D\\.\[0-9\]+:5:1\\):d\\\[\\(\\(integer\\(kind=8\\)\\) i \\+ -1\\) \\* 6\\\]\\)"  1 "original" } }
 
 ! { dg-final { scan-tree-dump-times "#pragma omp task affinity\\(iterator\\(integer\\(kind=4\\) i=1:5:1\\):a\\)\[^ \]" 1 "original" } }
 
 ! { dg-final { scan-tree-dump-times "#pragma omp task affinity\\(iterator\\(integer\\(kind=4\\) i=1:5:1\\):a\\) affinity\\(iterator\\(integer\\(kind=4\\) i=1:5:1\\):\\*x\\)"  1 "original" } }
 
-! { dg-final { scan-tree-dump-times "#pragma omp task affinity\\(iterator\\(integer\\(kind=4\\) k=7:4:-1, integer\\(kind=8\\) j=1:5:1\\):\\*\\(c_char \\*\\) &b\\\[\\(?\\(integer\\(kind=.\\).* \[jk\] \\+ .*\[kj\]\\) \\+ -1\\\]\\) affinity\\(iterator\\(integer\\(kind=4\\) k=7:4:-1, integer\\(kind=8\\) j=1:5:1\\):a\\) affinity\\(cc\\)" 1 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma omp task affinity\\(iterator\\(integer\\(kind=4\\) k=7:4:-1, integer\\(kind=8\\) j=1:5:1\\):b\\\[\\(?\\(integer\\(kind=.\\).* \[jk\] \\+ .*\[kj\]\\) \\+ -1\\\]\\) affinity\\(iterator\\(integer\\(kind=4\\) k=7:4:-1, integer\\(kind=8\\) j=1:5:1\\):a\\) affinity\\(cc\\)" 1 "original" } }

diff --git a/gcc/testsuite/gfortran.dg/gomp/affinity-clause-5.f90 b/gcc/testsuite/gfortran.dg/gomp/affinity-clause-5.f90
index 538b5a5..c23fee0 100644 (file)
--- a/gcc/testsuite/gfortran.dg/gomp/affinity-clause-5.f90
+++ b/gcc/testsuite/gfortran.dg/gomp/affinity-clause-5.f90
@@ -18,6 +18,6 @@ end
 
 ! { dg-final { scan-tree-dump-times "pragma omp task affinity\\(iterator\\)" 1 "original" } }
 
-! { dg-final { scan-tree-dump-times "#pragma omp task affinity\\(\\*\\(c_char \\*\\) &iterator\\\[2\\\]\\)" 1 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma omp task affinity\\(iterator\\\[2\\\]\\)" 1 "original" } }
 
-! { dg-final { scan-tree-dump-times "#pragma omp task affinity\\(iterator\\(integer\\(kind=4\\) i=1:10:1\\):\\*\\(c_char \\*\\) &iterator\\\[.* <?i>? \\+ -1\\\]\\)" 1 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma omp task affinity\\(iterator\\(integer\\(kind=4\\) i=1:10:1\\):iterator\\\[.* <?i>? \\+ -1\\\]\\)" 1 "original" } }

diff --git a/gcc/testsuite/gfortran.dg/gomp/defaultmap-4.f90 b/gcc/testsuite/gfortran.dg/gomp/defaultmap-4.f90
index 89bbe87..7b182b5 100644 (file)
--- a/gcc/testsuite/gfortran.dg/gomp/defaultmap-4.f90
+++ b/gcc/testsuite/gfortran.dg/gomp/defaultmap-4.f90
@@ -56,16 +56,18 @@ end
 ! { dg-final { scan-tree-dump-times "map\\(alloc:\\*aii \\\[len:" 1 "gimple" } }
 ! { dg-final { scan-tree-dump-times "map\\(alloc:aii \\\[pointer assign, bias: 0\\\]\\)" 2 "gimple" } }
 ! { dg-final { scan-tree-dump-times "map\\(alloc:arr \\\[len:" 1 "gimple" } }
-! { dg-final { scan-tree-dump-times "map\\(alloc:\\*\\(c_char \\*\\) aarr\\.data \\\[len:" 1 "gimple" } }
-! { dg-final { scan-tree-dump-times "map\\(alloc:\\*\\(c_char \\*\\) dtaarr\\.data \\\[len:" 1 "gimple" } }
-! { dg-final { scan-tree-dump-times "map\\(alloc:\\*\\(c_char \\*\\) dtparr\\.data \\\[len:" 1 "gimple" } }
-! { dg-final { scan-tree-dump-times "map\\(alloc:\\*\\(c_char \\*\\) parr\\.data \\\[len:" 1 "gimple" } }
-! { dg-final { scan-tree-dump-times "map\\(alloc:\\*\\(c_char \\*\\) str1aarr\\.data \\\[len:" 1 "gimple" } }
-! { dg-final { scan-tree-dump-times "map\\(alloc:\\*\\(c_char \\*\\) str1parr\\.data \\\[len:" 1 "gimple" } }
-! { dg-final { scan-tree-dump-times "map\\(alloc:\\*\\(c_char \\*\\) str5aarr\\.data \\\[len:" 1 "gimple" } }
-! { dg-final { scan-tree-dump-times "map\\(alloc:\\*\\(c_char \\*\\) str5parr\\.data \\\[len:" 1 "gimple" } }
-! { dg-final { scan-tree-dump-times "map\\(alloc:\\*\\(c_char \\*\\) strxaarr\\.data \\\[len:" 1 "gimple" } }
-! { dg-final { scan-tree-dump-times "map\\(alloc:\\*\\(c_char \\*\\) strxparr\\.data \\\[len:" 1 "gimple" } }
+
+! { dg-final { scan-tree-dump-times "map\\(alloc:\\*\\(character\\(kind=1\\)\\\[0:\\\]\\\[1:\\.strxparr\\\] \\*\\) strxparr\\.data \\\[len:" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "map\\(alloc:\\*\\(character\\(kind=1\\)\\\[0:\\\]\\\[1:\\.strxaarr\\\] \\* restrict\\) strxaarr\\.data \\\[len:" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "map\\(alloc:\\*\\(character\\(kind=1\\)\\\[0:\\\]\\\[1:5\\\] \\*\\) str5parr\\.data \\\[len:" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "map\\(alloc:\\*\\(character\\(kind=1\\)\\\[0:\\\]\\\[1:5\\\] \\* restrict\\) str5aarr\\.data \\\[len:" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "map\\(alloc:\\*\\(character\\(kind=1\\)\\\[0:\\\]\\\[1:1\\\] \\*\\) str1parr\\.data \\\[len:" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "map\\(alloc:\\*\\(character\\(kind=1\\)\\\[0:\\\]\\\[1:1\\\] \\* restrict\\) str1aarr\\.data \\\[len:" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "map\\(alloc:\\*\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) parr\\.data \\\[len:" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "map\\(alloc:\\*\\(struct t\\\[0:\\\] \\*\\) dtparr\\.data \\\[len:" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "map\\(alloc:\\*\\(struct t\\\[0:\\\] \\* restrict\\) dtaarr\\.data \\\[len:" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "map\\(alloc:\\*\\(integer\\(kind=4\\)\\\[0:\\\] \\* restrict\\) aarr\\.data \\\[len:" 1 "gimple" } }
+
 ! { dg-final { scan-tree-dump-times "map\\(alloc:\\(character\\(kind=1\\)\\\[0:\\\]\\\[1:1\\\] \\* restrict\\) str1aarr\\.data \\\[pointer assign, bias: 0\\\]\\)" 2 "gimple" } }
 ! { dg-final { scan-tree-dump-times "map\\(alloc:\\(character\\(kind=1\\)\\\[0:\\\]\\\[1:5\\\] \\* restrict\\) str5aarr\\.data \\\[pointer assign, bias: 0\\\]\\)" 2 "gimple" } }
 ! { dg-final { scan-tree-dump-times "map\\(alloc:\\(character\\(kind=1\\)\\\[0:\\\]\\\[1:\\.strxaarr\\\] \\* restrict\\) strxaarr\\.data \\\[pointer assign, bias: 0\\\]\\)" 2 "gimple" } }
@@ -103,21 +105,23 @@ end
 ! { dg-final { scan-tree-dump-times "map\\(always_pointer:\\(struct t\\\[0:\\\] \\*\\) dtparr\\.data \\\[pointer assign, bias: 0\\\]\\)" 2 "gimple" } }
 ! { dg-final { scan-tree-dump-times "map\\(to:aarr \\\[pointer set, len:" 2 "gimple" } }
 ! { dg-final { scan-tree-dump-times "map\\(to:arr \\\[len:" 1 "gimple" } }
-! { dg-final { scan-tree-dump-times "map\\(to:\\*\\(c_char \\*\\) dtparr\\.data \\\[len:" 1 "gimple" } }
-! { dg-final { scan-tree-dump-times "map\\(to:\\*\\(c_char \\*\\) parr\\.data \\\[len:" 1 "gimple" } }
-! { dg-final { scan-tree-dump-times "map\\(to:\\*\\(c_char \\*\\) str1parr\\.data \\\[len:" 1 "gimple" } }
-! { dg-final { scan-tree-dump-times "map\\(to:\\*\\(c_char \\*\\) str5parr\\.data \\\[len:" 1 "gimple" } }
-! { dg-final { scan-tree-dump-times "map\\(to:\\*\\(c_char \\*\\) strxparr\\.data \\\[len:" 1 "gimple" } }
+
+! { dg-final { scan-tree-dump-times "map\\(to:\\*\\(character\\(kind=1\\)\\\[0:\\\]\\\[1:\\.strxparr\\\] \\*\\) strxparr\\.data \\\[len:" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "map\\(to:\\*\\(character\\(kind=1\\)\\\[0:\\\]\\\[1:5\\\] \\*\\) str5parr\\.data \\\[len:" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "map\\(to:\\*\\(character\\(kind=1\\)\\\[0:\\\]\\\[1:1\\\] \\*\\) str1parr\\.data \\\[len:" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "map\\(to:\\*\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) parr\\.data \\\[len:" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "map\\(to:\\*\\(struct t\\\[0:\\\] \\*\\) dtparr\\.data \\\[len:" 1 "gimple" } }
+
 ! { dg-final { scan-tree-dump-times "map\\(to:dtaarr \\\[pointer set, len:" 2 "gimple" } }
 ! { dg-final { scan-tree-dump-times "map\\(to:dtarr \\\[len:" 1 "gimple" } }
 ! { dg-final { scan-tree-dump-times "map\\(to:dt \\\[len:" 1 "gimple" } }
 ! { dg-final { scan-tree-dump-times "map\\(to:dtparr \\\[pointer set, len:" 2 "gimple" } }
 ! { dg-final { scan-tree-dump-times "map\\(tofrom:\\*aii \\\[len:" 1 "gimple" } }
-! { dg-final { scan-tree-dump-times "map\\(tofrom:\\*\\(c_char \\*\\) aarr\\.data \\\[len:" 1 "gimple" } }
-! { dg-final { scan-tree-dump-times "map\\(tofrom:\\*\\(c_char \\*\\) dtaarr\\.data \\\[len:" 1 "gimple" } }
-! { dg-final { scan-tree-dump-times "map\\(tofrom:\\*\\(c_char \\*\\) str1aarr\\.data \\\[len:" 1 "gimple" } }
-! { dg-final { scan-tree-dump-times "map\\(tofrom:\\*\\(c_char \\*\\) str5aarr\\.data \\\[len:" 1 "gimple" } }
-! { dg-final { scan-tree-dump-times "map\\(tofrom:\\*\\(c_char \\*\\) strxaarr\\.data \\\[len:" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "map\\(tofrom:\\*\\(character\\(kind=1\\)\\\[0:\\\]\\\[1:\\.strxaarr\\\] \\* restrict\\) strxaarr\\.data \\\[len:" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "map\\(tofrom:\\*\\(character\\(kind=1\\)\\\[0:\\\]\\\[1:5\\\] \\* restrict\\) str5aarr\\.data \\\[len:" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "map\\(tofrom:\\*\\(character\\(kind=1\\)\\\[0:\\\]\\\[1:1\\\] \\* restrict\\) str1aarr\\.data \\\[len:" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "map\\(tofrom:\\*\\(struct t\\\[0:\\\] \\* restrict\\) dtaarr\\.data \\\[len:" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "map\\(tofrom:\\*\\(integer\\(kind=4\\)\\\[0:\\\] \\* restrict\\) aarr\\.data \\\[len:" 1 "gimple" } }
 ! { dg-final { scan-tree-dump-times "map\\(tofrom:\\*dta \\\[len:" 1 "gimple" } }
 ! { dg-final { scan-tree-dump-times "map\\(tofrom:\\*str1a \\\[len:" 1 "gimple" } }
 ! { dg-final { scan-tree-dump-times "map\\(tofrom:\\*str5a \\\[len:" 1 "gimple" } }

diff --git a/gcc/testsuite/gfortran.dg/gomp/defaultmap-5.f90 b/gcc/testsuite/gfortran.dg/gomp/defaultmap-5.f90
index d6b32dc..1391274 100644 (file)
--- a/gcc/testsuite/gfortran.dg/gomp/defaultmap-5.f90
+++ b/gcc/testsuite/gfortran.dg/gomp/defaultmap-5.f90
@@ -86,16 +86,16 @@ end
 ! { dg-final { scan-tree-dump-times "map\\(to:aarr \\\[pointer set, len:" 1 "gimple" } }
 ! { dg-final { scan-tree-dump-times "map\\(to:\\*aii \\\[len:" 1 "gimple" } }
 ! { dg-final { scan-tree-dump-times "map\\(to:arr \\\[len:" 1 "gimple" } }
-! { dg-final { scan-tree-dump-times "map\\(to:\\*\\(c_char \\*\\) aarr\\.data \\\[len:" 1 "gimple" } }
-! { dg-final { scan-tree-dump-times "map\\(to:\\*\\(c_char \\*\\) dtaarr\\.data \\\[len:" 1 "gimple" } }
-! { dg-final { scan-tree-dump-times "map\\(to:\\*\\(c_char \\*\\) dtparr\\.data \\\[len:" 1 "gimple" } }
-! { dg-final { scan-tree-dump-times "map\\(to:\\*\\(c_char \\*\\) parr\\.data \\\[len:" 1 "gimple" } }
-! { dg-final { scan-tree-dump-times "map\\(to:\\*\\(c_char \\*\\) str1aarr\\.data \\\[len:" 1 "gimple" } }
-! { dg-final { scan-tree-dump-times "map\\(to:\\*\\(c_char \\*\\) str1parr\\.data \\\[len:" 1 "gimple" } }
-! { dg-final { scan-tree-dump-times "map\\(to:\\*\\(c_char \\*\\) str5aarr\\.data \\\[len:" 1 "gimple" } }
-! { dg-final { scan-tree-dump-times "map\\(to:\\*\\(c_char \\*\\) str5parr\\.data \\\[len:" 1 "gimple" } }
-! { dg-final { scan-tree-dump-times "map\\(to:\\*\\(c_char \\*\\) strxaarr\\.data \\\[len:" 1 "gimple" } }
-! { dg-final { scan-tree-dump-times "map\\(to:\\*\\(c_char \\*\\) strxparr\\.data \\\[len:" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "map\\(to:\\*\\(character\\(kind=1\\)\\\[0:\\\]\\\[1:\\.strxparr\\\] \\*\\) strxparr\\.data \\\[len:" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "map\\(to:\\*\\(character\\(kind=1\\)\\\[0:\\\]\\\[1:\\.strxaarr\\\] \\* restrict\\) strxaarr\\.data \\\[len:" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "map\\(to:\\*\\(character\\(kind=1\\)\\\[0:\\\]\\\[1:5\\\] \\*\\) str5parr\\.data \\\[len:" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "map\\(to:\\*\\(character\\(kind=1\\)\\\[0:\\\]\\\[1:5\\\] \\* restrict\\) str5aarr\\.data \\\[len:" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "map\\(to:\\*\\(character\\(kind=1\\)\\\[0:\\\]\\\[1:1\\\] \\*\\) str1parr\\.data \\\[len:" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "map\\(to:\\*\\(character\\(kind=1\\)\\\[0:\\\]\\\[1:1\\\] \\* restrict\\) str1aarr\\.data \\\[len:" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "map\\(to:\\*\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) parr\\.data \\\[len:" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "map\\(to:\\*\\(struct t\\\[0:\\\] \\*\\) dtparr\\.data \\\[len:" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "map\\(to:\\*\\(struct t\\\[0:\\\] \\* restrict\\) dtaarr\\.data \\\[len:" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "map\\(to:\\*\\(integer\\(kind=4\\)\\\[0:\\\] \\* restrict\\) aarr\\.data \\\[len:" 1 "gimple" } }
 ! { dg-final { scan-tree-dump-times "map\\(to:dtaarr \\\[pointer set, len:" 1 "gimple" } }
 ! { dg-final { scan-tree-dump-times "map\\(to:\\*dta \\\[len:" 1 "gimple" } }
 ! { dg-final { scan-tree-dump-times "map\\(to:dtarr \\\[len:" 1 "gimple" } }
@@ -103,11 +103,11 @@ end
 ! { dg-final { scan-tree-dump-times "map\\(to:dtparr \\\[pointer set, len:" 2 "gimple" } }
 ! { dg-final { scan-tree-dump-times "map\\(to:\\*dtp \\\[len:" 1 "gimple" } }
 ! { dg-final { scan-tree-dump-times "map\\(tofrom:arr \\\[len:" 1 "gimple" } }
-! { dg-final { scan-tree-dump-times "map\\(tofrom:\\*\\(c_char \\*\\) dtparr\\.data \\\[len:" 1 "gimple" } }
-! { dg-final { scan-tree-dump-times "map\\(tofrom:\\*\\(c_char \\*\\) parr\\.data \\\[len:" 1 "gimple" } }
-! { dg-final { scan-tree-dump-times "map\\(tofrom:\\*\\(c_char \\*\\) str1parr\\.data \\\[len:" 1 "gimple" } }
-! { dg-final { scan-tree-dump-times "map\\(tofrom:\\*\\(c_char \\*\\) str5parr\\.data \\\[len:" 1 "gimple" } }
-! { dg-final { scan-tree-dump-times "map\\(tofrom:\\*\\(c_char \\*\\) strxparr\\.data \\\[len:" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "map\\(tofrom:\\*\\(character\\(kind=1\\)\\\[0:\\\]\\\[1:\\.strxparr\\\] \\*\\) strxparr\\.data \\\[len:" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "map\\(tofrom:\\*\\(character\\(kind=1\\)\\\[0:\\\]\\\[1:5\\\] \\*\\) str5parr\\.data \\\[len:" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "map\\(tofrom:\\*\\(character\\(kind=1\\)\\\[0:\\\]\\\[1:1\\\] \\*\\) str1parr\\.data \\\[len:" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "map\\(tofrom:\\*\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) parr\\.data \\\[len:" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "map\\(tofrom:\\*\\(struct t\\\[0:\\\] \\*\\) dtparr\\.data \\\[len:" 1 "gimple" } }
 ! { dg-final { scan-tree-dump-times "map\\(tofrom:dtarr \\\[len:" 1 "gimple" } }
 ! { dg-final { scan-tree-dump-times "map\\(tofrom:dt \\\[len:" 1 "gimple" } }
 ! { dg-final { scan-tree-dump-times "map\\(tofrom:\\*dtp \\\[len:" 1 "gimple" } }

diff --git a/gcc/testsuite/gfortran.dg/gomp/defaultmap-6.f90 b/gcc/testsuite/gfortran.dg/gomp/defaultmap-6.f90
index fabf771..9a81d0f 100644 (file)
--- a/gcc/testsuite/gfortran.dg/gomp/defaultmap-6.f90
+++ b/gcc/testsuite/gfortran.dg/gomp/defaultmap-6.f90
@@ -65,16 +65,16 @@ end
 ! { dg-final { scan-tree-dump-times "map\\(to:dtparr \\\[pointer set, len:" 1 "gimple" } }
 ! { dg-final { scan-tree-dump-times "map\\(tofrom:\\*aii \\\[len:" 1 "gimple" } }
 ! { dg-final { scan-tree-dump-times "map\\(tofrom:arr \\\[len:" 1 "gimple" } }
-! { dg-final { scan-tree-dump-times "map\\(tofrom:\\*\\(c_char \\*\\) aarr\\.data \\\[len:" 1 "gimple" } }
-! { dg-final { scan-tree-dump-times "map\\(tofrom:\\*\\(c_char \\*\\) dtaarr\\.data \\\[len:" 1 "gimple" } }
-! { dg-final { scan-tree-dump-times "map\\(tofrom:\\*\\(c_char \\*\\) dtparr\\.data \\\[len:" 1 "gimple" } }
-! { dg-final { scan-tree-dump-times "map\\(tofrom:\\*\\(c_char \\*\\) parr\\.data \\\[len:" 1 "gimple" } }
-! { dg-final { scan-tree-dump-times "map\\(tofrom:\\*\\(c_char \\*\\) str1aarr\\.data \\\[len:" 1 "gimple" } }
-! { dg-final { scan-tree-dump-times "map\\(tofrom:\\*\\(c_char \\*\\) str1parr\\.data \\\[len:" 1 "gimple" } }
-! { dg-final { scan-tree-dump-times "map\\(tofrom:\\*\\(c_char \\*\\) str5aarr\\.data \\\[len:" 1 "gimple" } }
-! { dg-final { scan-tree-dump-times "map\\(tofrom:\\*\\(c_char \\*\\) str5parr\\.data \\\[len:" 1 "gimple" } }
-! { dg-final { scan-tree-dump-times "map\\(tofrom:\\*\\(c_char \\*\\) strxaarr\\.data \\\[len:" 1 "gimple" } }
-! { dg-final { scan-tree-dump-times "map\\(tofrom:\\*\\(c_char \\*\\) strxparr\\.data \\\[len:" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "map\\(tofrom:\\*\\(character\\(kind=1\\)\\\[0:\\\]\\\[1:\\.strxparr\\\] \\*\\) strxparr\\.data \\\[len:" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "map\\(tofrom:\\*\\(character\\(kind=1\\)\\\[0:\\\]\\\[1:\\.strxaarr\\\] \\* restrict\\) strxaarr\\.data \\\[len:" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "map\\(tofrom:\\*\\(character\\(kind=1\\)\\\[0:\\\]\\\[1:5\\\] \\*\\) str5parr\\.data \\\[len:" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "map\\(tofrom:\\*\\(character\\(kind=1\\)\\\[0:\\\]\\\[1:5\\\] \\* restrict\\) str5aarr\\.data \\\[len:" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "map\\(tofrom:\\*\\(character\\(kind=1\\)\\\[0:\\\]\\\[1:1\\\] \\*\\) str1parr\\.data \\\[len:" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "map\\(tofrom:\\*\\(character\\(kind=1\\)\\\[0:\\\]\\\[1:1\\\] \\* restrict\\) str1aarr\\.data \\\[len:" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "map\\(tofrom:\\*\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) parr\\.data \\\[len:" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "map\\(tofrom:\\*\\(struct t\\\[0:\\\] \\*\\) dtparr\\.data \\\[len:" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "map\\(tofrom:\\*\\(struct t\\\[0:\\\] \\* restrict\\) dtaarr\\.data \\\[len:" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "map\\(tofrom:\\*\\(integer\\(kind=4\\)\\\[0:\\\] \\* restrict\\) aarr\\.data \\\[len:" 1 "gimple" } }
 ! { dg-final { scan-tree-dump-times "map\\(tofrom:\\*dta \\\[len:" 1 "gimple" } }
 ! { dg-final { scan-tree-dump-times "map\\(tofrom:dtarr \\\[len:" 1 "gimple" } }
 ! { dg-final { scan-tree-dump-times "map\\(tofrom:dt \\\[len:" 1 "gimple" } }

diff --git a/gcc/testsuite/gfortran.dg/gomp/map-3.f90 b/gcc/testsuite/gfortran.dg/gomp/map-3.f90
index bdd2890..2f0a792 100644 (file)
--- a/gcc/testsuite/gfortran.dg/gomp/map-3.f90
+++ b/gcc/testsuite/gfortran.dg/gomp/map-3.f90
@@ -34,5 +34,5 @@ end
 ! { dg-final { scan-tree-dump-times "#pragma omp target data use_device_addr\\(x\\)" 1 "original" } }
 ! { dg-final { scan-tree-dump-times "#pragma omp target data use_device_addr\\(x2\\)" 1 "original" } }
 ! { dg-final { scan-tree-dump-times "#pragma omp target exit data map\\(release:x\\)" 1 "original" } }
-! { dg-final { scan-tree-dump-times "#pragma omp target data map\\(tofrom:\\*\\(c_char \\*\\) y.data \\\[len: .*\\) map\\(to:y \\\[pointer set, len: .*\\) map\\(alloc:.*y.data \\\[pointer assign, bias: 0\\\]\\) use_device_addr\\(y\\)" 1 "original" } }
-! { dg-final { scan-tree-dump-times "#pragma omp target data map\\(tofrom:\\*\\(c_char \\*\\) z.data \\\[len: .*\\) map\\(to:z \\\[pointer set, len: .*\\) map\\(alloc:.*z.data \\\[pointer assign, bias: 0\\\]\\) use_device_addr\\(z\\)" 1 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma omp target data map\\(tofrom:\\*\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) y.data \\\[len: .*\\) map\\(to:y \\\[pointer set, len: .*\\) map\\(alloc:.*y.data \\\[pointer assign, bias: 0\\\]\\) use_device_addr\\(y\\)" 1 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma omp target data map\\(tofrom:\\*\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) z.data \\\[len: .*\\) map\\(to:z \\\[pointer set, len: .*\\) map\\(alloc:.*z.data \\\[pointer assign, bias: 0\\\]\\) use_device_addr\\(z\\)" 1 "original" } }

diff --git a/gcc/testsuite/gfortran.dg/gomp/pr78260-2.f90 b/gcc/testsuite/gfortran.dg/gomp/pr78260-2.f90
index c58ad93..f5d8885 100644 (file)
--- a/gcc/testsuite/gfortran.dg/gomp/pr78260-2.f90
+++ b/gcc/testsuite/gfortran.dg/gomp/pr78260-2.f90
@@ -48,10 +48,10 @@ contains
   end subroutine sub
 end module m
 
-! { dg-final { scan-tree-dump-times "#pragma omp target data map\\(tofrom:\\*\\(c_char \\*\\) arr.data \\\[len: D.\[0-9\]+ \\* 4\\\]\\) map\\(to:arr \\\[pointer set, len: ..\\\]\\) map\\(alloc:\\(integer\\(kind=4\\)\\\[0:\\\] \\* restrict\\) arr.data \\\[pointer assign, bias: 0\\\]\\)" 1 "original" } }
-! { dg-final { scan-tree-dump-times "#pragma omp target update to\\(\\*\\(c_char \\*\\) arr.data \\\[len: D.\[0-9\]+ \\* 4\\\]\\)" 1 "original" } }
-! { dg-final { scan-tree-dump-times "#pragma omp target data map\\(tofrom:\\*\\(c_char \\*\\) __result->data \\\[len: D.\[0-9\]+ \\* 4\\\]\\) map\\(to:\\*__result \\\[pointer set, len: ..\\\]\\) map\\(alloc:\\(integer\\(kind=4\\)\\\[0:\\\] \\* restrict\\) __result->data \\\[pointer assign, bias: 0\\\]\\) map\\(alloc:__result \\\[pointer assign, bias: 0\\\]\\)" 1 "original" } }
-! { dg-final { scan-tree-dump-times "#pragma omp target update to\\(\\*\\(c_char \\*\\) __result->data \\\[len: D.\[0-9\]+ \\* 4\\\]\\)" 1 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma omp target data map\\(tofrom:\\*\\(integer\\(kind=4\\)\\\[0:\\\] \\* restrict\\) arr.data \\\[len: D.\[0-9\]+ \\* 4\\\]\\) map\\(to:arr \\\[pointer set, len: ..\\\]\\) map\\(alloc:\\(integer\\(kind=4\\)\\\[0:\\\] \\* restrict\\) arr.data \\\[pointer assign, bias: 0\\\]\\)" 1 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma omp target update to\\(\\*\\(integer\\(kind=4\\)\\\[0:\\\] \\* restrict\\) arr.data \\\[len: D.\[0-9\]+ \\* 4\\\]\\)" 1 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma omp target data map\\(tofrom:\\*\\(integer\\(kind=4\\)\\\[0:\\\] \\* restrict\\) __result->data \\\[len: D.\[0-9\]+ \\* 4\\\]\\) map\\(to:\\*__result \\\[pointer set, len: ..\\\]\\) map\\(alloc:\\(integer\\(kind=4\\)\\\[0:\\\] \\* restrict\\) __result->data \\\[pointer assign, bias: 0\\\]\\) map\\(alloc:__result \\\[pointer assign, bias: 0\\\]\\)" 1 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma omp target update to\\(\\*\\(integer\\(kind=4\\)\\\[0:\\\] \\* restrict\\) __result->data \\\[len: D.\[0-9\]+ \\* 4\\\]\\)" 1 "original" } }
 ! { dg-final { scan-tree-dump-times "#pragma omp target data map\\(tofrom:\\*__result.0\\) map\\(alloc:__result.0 \\\[pointer assign, bias: 0\\\]\\)" 2 "original" } }
 ! { dg-final { scan-tree-dump-times "#pragma omp target update to\\(\\*__result.0\\)" 2 "original" } }
 ! { dg-final { scan-tree-dump-times "#pragma omp target data map\\(tofrom:__result_f1\\)" 1 "original" } }

diff --git a/gcc/testsuite/gfortran.dg/gomp/pr78260-3.f90 b/gcc/testsuite/gfortran.dg/gomp/pr78260-3.f90
index 4ca3e36..64851b3 100644 (file)
--- a/gcc/testsuite/gfortran.dg/gomp/pr78260-3.f90
+++ b/gcc/testsuite/gfortran.dg/gomp/pr78260-3.f90
@@ -70,5 +70,5 @@ end subroutine sub
 
 ! { dg-final { scan-tree-dump-times "#pragma omp task depend\\(inout:__result_f1\\)" 2 "original" } }
 ! { dg-final { scan-tree-dump-times "#pragma omp task depend\\(inout:\\*__result.0\\)" 4 "original" } }
-! { dg-final { scan-tree-dump-times "#pragma omp task depend\\(inout:\\*\\(c_char \\*\\) __result->data\\)" 2 "original" } }
-! { dg-final { scan-tree-dump-times "#pragma omp task depend\\(inout:\\*\\(c_char \\*\\) arr.data\\)" 2 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma omp task depend\\(inout:\\*\\(integer\\(kind=4\\)\\\[0:\\\] \\* restrict\\) __result->data\\)" 2 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma omp task depend\\(inout:\\*\\(integer\\(kind=4\\)\\\[0:\\\] \\* restrict\\) arr.data\\)" 2 "original" } }

diff --git a/libgomp/testsuite/libgomp.fortran/pr90030.f90 b/libgomp/testsuite/libgomp.fortran/pr90030.f90
new file mode 100644 (file)
index 0000000..8c2432c
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/pr90030.f90
@@ -0,0 +1,3 @@
+! { dg-do run }
+
+include '../libgomp.oacc-fortran/pr90030.f90'

diff --git a/libgomp/testsuite/libgomp.oacc-fortran/pr90030.f90 b/libgomp/testsuite/libgomp.oacc-fortran/pr90030.f90
new file mode 100644 (file)
index 0000000..bbfcff3
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/pr90030.f90
@@ -0,0 +1,29 @@
+! PR90030.
+! Test if the array data associated with c is properly aligned
+! on the accelerator.  If it is not, this program will crash.
+
+! This is also included from '../libgomp.fortran/pr90030.f90'.
+
+! { dg-do run }
+
+program routine_align_main
+  implicit none
+  integer :: i, n
+  real*8, dimension(:), allocatable :: c
+
+  n = 10
+
+  allocate (c(n))
+
+  !$omp target map(to: n) map(from: c(1:n))
+  !$acc parallel copyin(n) copyout(c(1:n))
+  do i = 1, n
+     c(i) = i
+  enddo
+  !$acc end parallel
+  !$omp end target
+
+  do i = 1, n
+     if (c(i) .ne. i) stop i
+  enddo
+end program routine_align_main