From: Chung-Lin Tang Date: Thu, 2 Dec 2021 10:24:03 +0000 (+0800) Subject: fortran: OpenMP/OpenACC array mapping alignment fix (PR90030) X-Git-Url: http://review.tizen.org/git/?a=commitdiff_plain;h=1ac7a8c9e4798d352eb8c64905dd38086af4e1cd;p=test_jj.git 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 --- 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 --- 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 --- 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\\\[.* ? \\+ -1\\\]\\) affinity\\(iterator\\(integer\\(kind=4\\) jj=2:5:2, integer\\(kind=4\\) i=D\\.\[0-9\]+:5:1\\):\\*\\(c_char \\*\\) &d\\\[\\(.*jj \\* 5 \\+ .* ?\\) \\+ -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\\\[.* ? \\+ -1\\\]\\) affinity\\(iterator\\(integer\\(kind=4\\) jj=2:5:2, integer\\(kind=4\\) i=D\\.\[0-9\]+:5:1\\):d\\\[\\(.*jj \\* 5 \\+ .* ?\\) \\+ -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\\\[\\(.* ? \\+ -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\\\[\\(.* ? \\+ -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 --- 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\\\[.* ? \\+ -1\\\]\\)" 1 "original" } } +! { dg-final { scan-tree-dump-times "#pragma omp task affinity\\(iterator\\(integer\\(kind=4\\) i=1:10:1\\):iterator\\\[.* ? \\+ -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 --- 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 --- 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 --- 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 --- 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 --- 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 --- 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 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 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