public inbox for gcc-cvs@sourceware.org
help / color / mirror / Atom feed
* [gcc/devel/omp/gcc-11] fortran: OpenMP/OpenACC array mapping alignment fix (PR90030)
@ 2021-12-08 14:38 Chung-Lin Tang
0 siblings, 0 replies; only message in thread
From: Chung-Lin Tang @ 2021-12-08 14:38 UTC (permalink / raw)
To: gcc-cvs
https://gcc.gnu.org/g:8f7d25feef166299fb9da89281ec857c565782dd
commit 8f7d25feef166299fb9da89281ec857c565782dd
Author: Chung-Lin Tang <cltang@codesourcery.com>
Date: Thu Dec 2 18:24:03 2021 +0800
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:
---
gcc/fortran/trans-openmp.c | 19 ++++------
gcc/testsuite/gfortran.dg/goacc/finalize-1.f | 8 ++--
.../gfortran.dg/gomp/affinity-clause-1.f90 | 6 +--
.../gfortran.dg/gomp/affinity-clause-5.f90 | 4 +-
gcc/testsuite/gfortran.dg/gomp/defaultmap-4.f90 | 44 ++++++++++++----------
gcc/testsuite/gfortran.dg/gomp/defaultmap-5.f90 | 30 +++++++--------
gcc/testsuite/gfortran.dg/gomp/defaultmap-6.f90 | 20 +++++-----
gcc/testsuite/gfortran.dg/gomp/map-3.f90 | 4 +-
gcc/testsuite/gfortran.dg/gomp/pr78260-2.f90 | 8 ++--
gcc/testsuite/gfortran.dg/gomp/pr78260-3.f90 | 4 +-
libgomp/testsuite/libgomp.fortran/pr90030.f90 | 3 ++
libgomp/testsuite/libgomp.oacc-fortran/pr90030.f90 | 29 ++++++++++++++
12 files changed, 105 insertions(+), 74 deletions(-)
diff --git a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c
index 2611d9b6ba8..ea16e27ff39 100644
--- 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 a7788580819..233b4cacf84 100644
--- 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 13bdd36d0b4..08c7740cf0d 100644
--- 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 538b5a54191..c23fee06646 100644
--- 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 89bbe87e570..7b182b52021 100644
--- 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 d6b32dc90bc..1391274be31 100644
--- 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 fabf771f9c7..9a81d0f1532 100644
--- 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 bdd2890b277..2f0a7928819 100644
--- 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 c58ad93471c..f5d888592b9 100644
--- 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 4ca3e361a59..64851b34249 100644
--- 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
index 00000000000..8c2432cb178
--- /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
index 00000000000..bbfcff3a869
--- /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
^ permalink raw reply [flat|nested] only message in thread
only message in thread, other threads:[~2021-12-08 14:38 UTC | newest]
Thread overview: (only message) (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2021-12-08 14:38 [gcc/devel/omp/gcc-11] fortran: OpenMP/OpenACC array mapping alignment fix (PR90030) Chung-Lin Tang
This is a public inbox, see mirroring instructions
for how to clone and mirror all data and code used for this inbox;
as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).