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.

(cherry picked from commit 1ac7a8c9e4798d352eb8c64905dd38086af4e1cd)

diff --git a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c
index 2611d9b6ba81b86a29a65da4a1f0c8c0bc037c85..ea16e27ff39e78874005d7e8a521b128f7f9dda7 100644 (file)
--- a/gcc/fortran/trans-openmp.c
+++ b/gcc/fortran/trans-openmp.c
@@ -1580,7 +1580,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);
@@ -2396,7 +2396,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);
 
@@ -2867,8 +2867,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))
@@ -2891,8 +2890,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)
@@ -3142,8 +3140,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,
@@ -3580,8 +3577,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)
@@ -3631,8 +3627,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 a77885808199357530440abfd877261ae96b25e1..233b4cacf8485ff90003d3692c7f37a9adfb531a 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_enter_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_enter_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_enter_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_enter_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 13bdd36d0b4d37365d17da1021625c96235a276d..08c7740cf0db053d4348f2ba6033447d4fb10d56 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 538b5a54191fa81cbe878b707a7c123b562802d1..c23fee066462f5f8c81338f7dc667a63d5da8868 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 89bbe87e570feb50a4f85dc7d277a7f1c4c90d0c..7b182b52021231a49d7c88733e75260dd93e2c1f 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 d6b32dc90bcc3f8e43e0db89ad8b46ce15f40652..1391274be31cf672a6cd7a60cf44c24d42fd0f2a 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 fabf771f9c7e91e636215f13843f07a67c2042d3..9a81d0f15324ca381a68561b71d9eb7a4d884dea 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 bdd2890b277f6257649bb15c4ca08dc685d59888..2f0a79288197534985c9074ed0e21634927d2afe 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 c58ad93471c2b60588d6faefa803e9879c80465f..f5d888592b95f35a7f04cbe03d5b4e3ee5245a50 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 4ca3e361a59b3ce8875ea510e0e142e330ac117d..64851b34249465fb888365a73a963dfaedc79ccb 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