[gcc/devel/omp/gcc-16] Add OpenACC Fortran support for deviceptr and variable in common blocks
Sandra Loosemore via Gcc-cvs <[email protected]>
| Newsgroups | gmane.comp.gcc.cvs |
|---|---|
| Message-ID | <[email protected]> |
https://gcc.gnu.org/g:55b1181c8083561fe84709b751edac6c9eba28e7 commit 55b1181c8083561fe84709b751edac6c9eba28e7 Author: Cesar Philippidis <[email protected]> Date: Wed Apr 23 03:01:28 2025 +0000 Add OpenACC Fortran support for deviceptr and variable in common blocks This is a merge of these OG14 commits: 987bb4c25b4076eb54f043644bdf9988378be90d 9e8395708c0027ad1de871bae870c4b0185a74fd 2adb0ec35cd47b34d47c961f6ae46089e3e02cbc 4d29174a9602e6ea783ba0e9a7b1e38fb6913db5 gcc/fortran/ChangeLog * openmp.cc (gfc_match_omp_map_clause): Re-write handling of the deviceptr clause. Add new common_blocks argument. Propagate it to gfc_match_omp_variable_list. (gfc_match_omp_clauses): Update calls to gfc_match_omp_map_clauses. (resolve_positive_int_expr): Promote the warning to an error. (check_array_not_assumed): Remove pointer check. (resolve_oacc_nested_loops): Error on do concurrent loops. * trans-openmp.cc (gfc_omp_finish_clause): Don't create pointer data mappings for deviceptr clauses. (gfc_trans_omp_clauses): Likewise. gcc/ChangeLog * gimplify.cc (enum gimplify_omp_var_data): Add GOVD_DEVICETPR. (oacc_default_clause): Privatize fortran common blocks. (omp_notice_variable): Add GOVD_DEVICEPTR attribute when appropriate. Defer the expansion of DECL_VALUE_EXPR for common block decls. (gimplify_scan_omp_clauses): Add GOVD_DEVICEPTR attribute when appropriate. (gimplify_adjust_omp_clauses_1): Set GOMP_MAP_FORCE_DEVICEPTR for implicit deviceptr mappings. gcc/testsuite/ChangeLog * c-c++-common/goacc/deviceptr-4.c: Update. * gfortran.dg/goacc/loop-2-kernels-tile.f95: Update. * gfortran.dg/goacc/loop-2-parallel-tile.f95: Update. * gfortran.dg/goacc/loop-2-serial-tile.f95: Update. * gfortran.dg/goacc/sie.f95: Update. * gfortran.dg/goacc/tile-1.f90: Update. * gfortran.dg/gomp/num-teams-2.f90: Update. * gfortran.dg/gomp/pr67500.f90: Update. * gfortran.dg/gomp/pr77516.f90: Update. libgomp/ChangeLog * oacc-parallel.c (GOACC_parallel_keyed): Handle Fortran deviceptr clause. (GOACC_data_start): Likewise. * testsuite/libgomp.oacc-fortran/deviceptr-1.f90: New test. Co-Authored-By: James Norris <[email protected]> Co-Authored-By: Julian Brown <[email protected]> Co-Authored-By: Tobias Burnus <[email protected]> Co-Authored-By: Thomas Schwinge <[email protected]> Diff: --- gcc/fortran/openmp.cc | 5 +- gcc/fortran/trans-openmp.cc | 9 + gcc/gimplify.cc | 12 +- gcc/testsuite/c-c++-common/goacc/deviceptr-4.c | 2 +- .../gfortran.dg/goacc/loop-2-kernels-tile.f95 | 4 +- .../gfortran.dg/goacc/loop-2-parallel-tile.f95 | 4 +- .../gfortran.dg/goacc/loop-2-serial-tile.f95 | 4 +- gcc/testsuite/gfortran.dg/goacc/sie.f95 | 36 ++-- gcc/testsuite/gfortran.dg/goacc/tile-1.f90 | 16 +- gcc/testsuite/gfortran.dg/gomp/num-teams-2.f90 | 12 +- gcc/testsuite/gfortran.dg/gomp/pr67500.f90 | 8 +- gcc/testsuite/gfortran.dg/gomp/pr77516.f90 | 2 +- libgomp/oacc-parallel.c | 2 + .../testsuite/libgomp.oacc-fortran/deviceptr-1.f90 | 197 +++++++++++++++++++++ 14 files changed, 265 insertions(+), 48 deletions(-) diff --git a/gcc/fortran/openmp.cc b/gcc/fortran/openmp.cc index c4289abaf851..62f3f44a9f7f 100644 --- a/gcc/fortran/openmp.cc +++ b/gcc/fortran/openmp.cc @@ -8581,9 +8581,8 @@ resolve_positive_int_expr (gfc_expr *expr, const char *clause) if (expr->expr_type == EXPR_CONSTANT && expr->ts.type == BT_INTEGER && mpz_sgn (expr->value.integer) <= 0) - gfc_warning ((flag_openmp || flag_openmp_simd) ? OPT_Wopenmp : 0, - "INTEGER expression of %s clause at %L must be positive", - clause, &expr->where); + gfc_error ("INTEGER expression of %s clause at %L must be positive", + clause, &expr->where); } static void diff --git a/gcc/fortran/trans-openmp.cc b/gcc/fortran/trans-openmp.cc index 7258d29455eb..ab6f751a79e6 100644 --- a/gcc/fortran/trans-openmp.cc +++ b/gcc/fortran/trans-openmp.cc @@ -1662,6 +1662,9 @@ gfc_omp_finish_clause (tree c, gimple_seq *pre_p, bool openacc) return; } + if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FORCE_DEVICEPTR) + return; + tree c2 = NULL_TREE, c3 = NULL_TREE, c4 = NULL_TREE; tree present = gfc_omp_check_optional_argument (decl, true); tree orig_decl = NULL_TREE; @@ -4674,6 +4677,12 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, OMP_CLAUSE_SIZE (node2) = size_int (0); goto finalize_map_clause; } + else if (POINTER_TYPE_P (type) + && n->u.map.op == OMP_MAP_FORCE_DEVICEPTR) + { + OMP_CLAUSE_DECL (node) = decl; + goto finalize_map_clause; + } else if (POINTER_TYPE_P (type) && (gfc_omp_privatize_by_reference (decl) || GFC_DECL_GET_SCALAR_POINTER (decl) diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc index 105bbb7b759f..8ad882e99df7 100644 --- a/gcc/gimplify.cc +++ b/gcc/gimplify.cc @@ -188,6 +188,9 @@ enum gimplify_omp_var_data /* Flag to indicate this an allocator for the uses_allocators clause. */ GOVD_USES_ALLOCATORS_ALLOCATOR = 0x8000000, + /* Flag for OpenACC deviceptrs. */ + GOVD_DEVICEPTR = 0x1000000, + GOVD_DATA_SHARE_CLASS = (GOVD_SHARED | GOVD_PRIVATE | GOVD_FIRSTPRIVATE | GOVD_LASTPRIVATE | GOVD_REDUCTION | GOVD_LINEAR | GOVD_LOCAL) @@ -9467,6 +9470,7 @@ omp_notice_variable (struct gimplify_omp_ctx *ctx, tree decl, bool in_code) error ("variable %qE declared in enclosing " "%<host_data%> region", DECL_NAME (decl)); nflags |= GOVD_MAP; + nflags |= (n2->value & GOVD_DEVICEPTR); if (octx->region_type == ORT_ACC_DATA && (n2->value & GOVD_MAP_0LEN_ARRAY)) nflags |= GOVD_MAP_0LEN_ARRAY; @@ -14668,6 +14672,8 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_PRESENT_TO || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_PRESENT_TOFROM) flags |= GOVD_MAP_ALWAYS_TO; + else if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FORCE_DEVICEPTR) + flags |= GOVD_DEVICEPTR; goto do_add; @@ -15732,7 +15738,8 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n, void *data) | GOVD_MAP_FORCE | GOVD_MAP_FORCE_PRESENT | GOVD_MAP_ALLOC_ONLY - | GOVD_MAP_FROM_ONLY)) + | GOVD_MAP_FROM_ONLY + | GOVD_DEVICEPTR)) { case 0: kind = GOMP_MAP_TOFROM; @@ -15758,6 +15765,9 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n, void *data) case GOVD_MAP_FORCE_PRESENT | GOVD_MAP_ALLOC_ONLY: kind = GOMP_MAP_FORCE_PRESENT; break; + case GOVD_DEVICEPTR: + kind = GOMP_MAP_FORCE_DEVICEPTR; + break; default: gcc_unreachable (); } diff --git a/gcc/testsuite/c-c++-common/goacc/deviceptr-4.c b/gcc/testsuite/c-c++-common/goacc/deviceptr-4.c index db1b91633a6f..79a51620db94 100644 --- a/gcc/testsuite/c-c++-common/goacc/deviceptr-4.c +++ b/gcc/testsuite/c-c++-common/goacc/deviceptr-4.c @@ -8,4 +8,4 @@ subr (int *a) a[0] += 1.0; } -/* { dg-final { scan-tree-dump-times "#pragma omp target oacc_parallel.*map\\(tofrom:a" 1 "gimple" } } */ +/* { dg-final { scan-tree-dump-times "#pragma omp target oacc_parallel.*map\\(force_deviceptr:a" 1 "gimple" } } */ diff --git a/gcc/testsuite/gfortran.dg/goacc/loop-2-kernels-tile.f95 b/gcc/testsuite/gfortran.dg/goacc/loop-2-kernels-tile.f95 index afc8a278cacd..65425159a2c9 100644 --- a/gcc/testsuite/gfortran.dg/goacc/loop-2-kernels-tile.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/loop-2-kernels-tile.f95 @@ -29,7 +29,7 @@ program test DO j = 1,10 ENDDO ENDDO - !$acc loop tile(-1) ! { dg-warning "must be positive" } + !$acc loop tile(-1) ! { dg-error "must be positive" } do i = 1,10 enddo !$acc loop tile(i) ! { dg-error "constant expression" } @@ -82,7 +82,7 @@ program test DO j = 1,10 ENDDO ENDDO - !$acc kernels loop tile(-1) ! { dg-warning "must be positive" } + !$acc kernels loop tile(-1) ! { dg-error "must be positive" } do i = 1,10 enddo !$acc kernels loop tile(i) ! { dg-error "constant expression" } diff --git a/gcc/testsuite/gfortran.dg/goacc/loop-2-parallel-tile.f95 b/gcc/testsuite/gfortran.dg/goacc/loop-2-parallel-tile.f95 index 4bfca748f75e..dae8f6674865 100644 --- a/gcc/testsuite/gfortran.dg/goacc/loop-2-parallel-tile.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/loop-2-parallel-tile.f95 @@ -20,7 +20,7 @@ program test DO j = 1,10 ENDDO ENDDO - !$acc loop tile(-1) ! { dg-warning "must be positive" } + !$acc loop tile(-1) ! { dg-error "must be positive" } do i = 1,10 enddo !$acc loop tile(i) ! { dg-error "constant expression" } @@ -73,7 +73,7 @@ program test DO j = 1,10 ENDDO ENDDO - !$acc parallel loop tile(-1) ! { dg-warning "must be positive" } + !$acc parallel loop tile(-1) ! { dg-error "must be positive" } do i = 1,10 enddo !$acc parallel loop tile(i) ! { dg-error "constant expression" } diff --git a/gcc/testsuite/gfortran.dg/goacc/loop-2-serial-tile.f95 b/gcc/testsuite/gfortran.dg/goacc/loop-2-serial-tile.f95 index 79d8d1b9d477..2e0a7a0d9daf 100644 --- a/gcc/testsuite/gfortran.dg/goacc/loop-2-serial-tile.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/loop-2-serial-tile.f95 @@ -20,7 +20,7 @@ program test DO j = 1,10 ENDDO ENDDO - !$acc loop tile(-1) ! { dg-warning "must be positive" } + !$acc loop tile(-1) ! { dg-error "must be positive" } do i = 1,10 enddo !$acc loop tile(i) ! { dg-error "constant expression" } @@ -73,7 +73,7 @@ program test DO j = 1,10 ENDDO ENDDO - !$acc serial loop tile(-1) ! { dg-warning "must be positive" } + !$acc serial loop tile(-1) ! { dg-error "must be positive" } do i = 1,10 enddo !$acc serial loop tile(i) ! { dg-error "constant expression" } diff --git a/gcc/testsuite/gfortran.dg/goacc/sie.f95 b/gcc/testsuite/gfortran.dg/goacc/sie.f95 index 78b4601004c7..c424b4cd1f18 100644 --- a/gcc/testsuite/gfortran.dg/goacc/sie.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/sie.f95 @@ -107,10 +107,10 @@ program test !$acc parallel num_gangs(i+1) !$acc end parallel - !$acc parallel num_gangs(-1) ! { dg-warning "must be positive" } + !$acc parallel num_gangs(-1) ! { dg-error "must be positive" } !$acc end parallel - !$acc parallel num_gangs(0) ! { dg-warning "must be positive" } + !$acc parallel num_gangs(0) ! { dg-error "must be positive" } !$acc end parallel !$acc parallel num_gangs() ! { dg-error "Invalid character in name" } @@ -135,10 +135,10 @@ program test !$acc kernels num_gangs(i+1) !$acc end kernels - !$acc kernels num_gangs(-1) ! { dg-warning "must be positive" } + !$acc kernels num_gangs(-1) ! { dg-error "must be positive" } !$acc end kernels - !$acc kernels num_gangs(0) ! { dg-warning "must be positive" } + !$acc kernels num_gangs(0) ! { dg-error "must be positive" } !$acc end kernels !$acc kernels num_gangs() ! { dg-error "Invalid character in name" } @@ -164,10 +164,10 @@ program test !$acc parallel num_workers(i+1) !$acc end parallel - !$acc parallel num_workers(-1) ! { dg-warning "must be positive" } + !$acc parallel num_workers(-1) ! { dg-error "must be positive" } !$acc end parallel - !$acc parallel num_workers(0) ! { dg-warning "must be positive" } + !$acc parallel num_workers(0) ! { dg-error "must be positive" } !$acc end parallel !$acc parallel num_workers() ! { dg-error "Invalid expression after 'num_workers\\('" } @@ -192,10 +192,10 @@ program test !$acc kernels num_workers(i+1) !$acc end kernels - !$acc kernels num_workers(-1) ! { dg-warning "must be positive" } + !$acc kernels num_workers(-1) ! { dg-error "must be positive" } !$acc end kernels - !$acc kernels num_workers(0) ! { dg-warning "must be positive" } + !$acc kernels num_workers(0) ! { dg-error "must be positive" } !$acc end kernels !$acc kernels num_workers() ! { dg-error "Invalid expression after 'num_workers\\('" } @@ -221,10 +221,10 @@ program test !$acc parallel vector_length(i+1) !$acc end parallel - !$acc parallel vector_length(-1) ! { dg-warning "must be positive" } + !$acc parallel vector_length(-1) ! { dg-error "must be positive" } !$acc end parallel - !$acc parallel vector_length(0) ! { dg-warning "must be positive" } + !$acc parallel vector_length(0) ! { dg-error "must be positive" } !$acc end parallel !$acc parallel vector_length() ! { dg-error "Invalid expression after 'vector_length\\('" } @@ -249,10 +249,10 @@ program test !$acc kernels vector_length(i+1) !$acc end kernels - !$acc kernels vector_length(-1) ! { dg-warning "must be positive" } + !$acc kernels vector_length(-1) ! { dg-error "must be positive" } !$acc end kernels - !$acc kernels vector_length(0) ! { dg-warning "must be positive" } + !$acc kernels vector_length(0) ! { dg-error "must be positive" } !$acc end kernels !$acc kernels vector_length() ! { dg-error "Invalid expression after 'vector_length\\('" } @@ -279,10 +279,10 @@ program test !$acc loop gang(i+1) do i = 1,10 enddo - !$acc loop gang(-1) ! { dg-warning "must be positive" } + !$acc loop gang(-1) ! { dg-error "must be positive" } do i = 1,10 enddo - !$acc loop gang(0) ! { dg-warning "must be positive" } + !$acc loop gang(0) ! { dg-error "must be positive" } do i = 1,10 enddo !$acc loop gang() ! { dg-error "Invalid character in name" } @@ -311,10 +311,10 @@ program test !$acc loop worker(i+1) do i = 1,10 enddo - !$acc loop worker(-1) ! { dg-warning "must be positive" } + !$acc loop worker(-1) ! { dg-error "must be positive" } do i = 1,10 enddo - !$acc loop worker(0) ! { dg-warning "must be positive" } + !$acc loop worker(0) ! { dg-error "must be positive" } do i = 1,10 enddo !$acc loop worker() ! { dg-error "Invalid character in name" } @@ -343,10 +343,10 @@ program test !$acc loop vector(i+1) do i = 1,10 enddo - !$acc loop vector(-1) ! { dg-warning "must be positive" } + !$acc loop vector(-1) ! { dg-error "must be positive" } do i = 1,10 enddo - !$acc loop vector(0) ! { dg-warning "must be positive" } + !$acc loop vector(0) ! { dg-error "must be positive" } do i = 1,10 enddo !$acc loop vector() ! { dg-error "Invalid character in name" } diff --git a/gcc/testsuite/gfortran.dg/goacc/tile-1.f90 b/gcc/testsuite/gfortran.dg/goacc/tile-1.f90 index f609b127df98..9ef752110877 100644 --- a/gcc/testsuite/gfortran.dg/goacc/tile-1.f90 +++ b/gcc/testsuite/gfortran.dg/goacc/tile-1.f90 @@ -44,17 +44,17 @@ subroutine parloop do i = 1, n end do - !$acc parallel loop tile(-3) ! { dg-warning "must be positive" } + !$acc parallel loop tile(-3) ! { dg-error "must be positive" } do i = 1, n end do - !$acc parallel loop tile(10, -3) ! { dg-warning "must be positive" } + !$acc parallel loop tile(10, -3) ! { dg-error "must be positive" } do i = 1, n do j = 1, n end do end do - !$acc parallel loop tile(-100, 10, 5) ! { dg-warning "must be positive" } + !$acc parallel loop tile(-100, 10, 5) ! { dg-error "must be positive" } do i = 1, n do j = 1, n do k = 1, n @@ -114,7 +114,7 @@ subroutine par end do end do - !$acc loop tile(-2) ! { dg-warning "must be positive" } + !$acc loop tile(-2) ! { dg-error "must be positive" } do i = 1, n end do @@ -195,7 +195,7 @@ subroutine kern end do end do - !$acc loop tile(-2) ! { dg-warning "must be positive" } + !$acc loop tile(-2) ! { dg-error "must be positive" } do i = 1, n end do @@ -295,17 +295,17 @@ subroutine kernsloop do i = 1, n end do - !$acc kernels loop tile(-3) ! { dg-warning "must be positive" } + !$acc kernels loop tile(-3) ! { dg-error "must be positive" } do i = 1, n end do - !$acc kernels loop tile(10, -3) ! { dg-warning "must be positive" } + !$acc kernels loop tile(10, -3) ! { dg-error "must be positive" } do i = 1, n do j = 1, n end do end do - !$acc kernels loop tile(-100, 10, 5) ! { dg-warning "must be positive" } + !$acc kernels loop tile(-100, 10, 5) ! { dg-error "must be positive" } do i = 1, n do j = 1, n do k = 1, n diff --git a/gcc/testsuite/gfortran.dg/gomp/num-teams-2.f90 b/gcc/testsuite/gfortran.dg/gomp/num-teams-2.f90 index e7814a11a5a4..f3031481d4a1 100644 --- a/gcc/testsuite/gfortran.dg/gomp/num-teams-2.f90 +++ b/gcc/testsuite/gfortran.dg/gomp/num-teams-2.f90 @@ -9,13 +9,13 @@ subroutine foo (i) !$omp teams num_teams (6 : 4) ! { dg-warning "NUM_TEAMS lower bound at .1. larger than upper bound at .2." } !$omp end teams - !$omp teams num_teams (-7) ! { dg-warning "INTEGER expression of NUM_TEAMS clause at .1. must be positive" } + !$omp teams num_teams (-7) ! { dg-error "INTEGER expression of NUM_TEAMS clause at .1. must be positive" } !$omp end teams - !$omp teams num_teams (i : -7) ! { dg-warning "INTEGER expression of NUM_TEAMS clause at .1. must be positive" } + !$omp teams num_teams (i : -7) ! { dg-error "INTEGER expression of NUM_TEAMS clause at .1. must be positive" } !$omp end teams - !$omp teams num_teams (-7 : 8) ! { dg-warning "INTEGER expression of NUM_TEAMS clause at .1. must be positive" } + !$omp teams num_teams (-7 : 8) ! { dg-error "INTEGER expression of NUM_TEAMS clause at .1. must be positive" } !$omp end teams end @@ -25,13 +25,13 @@ subroutine bar (i) !$omp target teams num_teams (6 : 4) ! { dg-warning "NUM_TEAMS lower bound at .1. larger than upper bound at .2." } !$omp end target teams - !$omp target teams num_teams (-7) ! { dg-warning "INTEGER expression of NUM_TEAMS clause at .1. must be positive" } + !$omp target teams num_teams (-7) ! { dg-error "INTEGER expression of NUM_TEAMS clause at .1. must be positive" } !$omp end target teams - !$omp target teams num_teams (i : -7) ! { dg-warning "INTEGER expression of NUM_TEAMS clause at .1. must be positive" } + !$omp target teams num_teams (i : -7) ! { dg-error "INTEGER expression of NUM_TEAMS clause at .1. must be positive" } !$omp end target teams - !$omp target teams num_teams (-7 : 8) ! { dg-warning "INTEGER expression of NUM_TEAMS clause at .1. must be positive" } + !$omp target teams num_teams (-7 : 8) ! { dg-error "INTEGER expression of NUM_TEAMS clause at .1. must be positive" } !$omp end target teams end end module diff --git a/gcc/testsuite/gfortran.dg/gomp/pr67500.f90 b/gcc/testsuite/gfortran.dg/gomp/pr67500.f90 index 1cecdc485781..11ed69f10a7b 100644 --- a/gcc/testsuite/gfortran.dg/gomp/pr67500.f90 +++ b/gcc/testsuite/gfortran.dg/gomp/pr67500.f90 @@ -10,11 +10,11 @@ subroutine f2 end subroutine f3 (i) - !$omp declare simd simdlen(-2) ! { dg-warning "INTEGER expression of SIMDLEN clause at .1. must be positive" } + !$omp declare simd simdlen(-2) ! { dg-error "INTEGER expression of SIMDLEN clause at .1. must be positive" } end subroutine subroutine f4 - !$omp declare simd simdlen(0) ! { dg-warning "INTEGER expression of SIMDLEN clause at .1. must be positive" } + !$omp declare simd simdlen(0) ! { dg-error "INTEGER expression of SIMDLEN clause at .1. must be positive" } end subroutine foo(p, d, n) @@ -31,11 +31,11 @@ subroutine foo(p, d, n) do i = 1, 16 end do - !$omp simd safelen(-2) ! { dg-warning "INTEGER expression of SAFELEN clause at .1. must be positive" } + !$omp simd safelen(-2) ! { dg-error "INTEGER expression of SAFELEN clause at .1. must be positive" } do i = 1, 16 end do - !$omp simd safelen(0) ! { dg-warning "INTEGER expression of SAFELEN clause at .1. must be positive" } + !$omp simd safelen(0) ! { dg-error "INTEGER expression of SAFELEN clause at .1. must be positive" } do i = 1, 16 end do diff --git a/gcc/testsuite/gfortran.dg/gomp/pr77516.f90 b/gcc/testsuite/gfortran.dg/gomp/pr77516.f90 index 9c0a95b9f79a..3ac3f5562d0a 100644 --- a/gcc/testsuite/gfortran.dg/gomp/pr77516.f90 +++ b/gcc/testsuite/gfortran.dg/gomp/pr77516.f90 @@ -4,7 +4,7 @@ program pr77516 integer :: i, x x = 0 -!$omp simd safelen(0) reduction(+:x) ! { dg-warning "must be positive" } +!$omp simd safelen(0) reduction(+:x) ! { dg-error "must be positive" } do i = 1, 8 x = x + 1 end do diff --git a/libgomp/oacc-parallel.c b/libgomp/oacc-parallel.c index fd4ab245f977..bf2b989bdbf4 100644 --- a/libgomp/oacc-parallel.c +++ b/libgomp/oacc-parallel.c @@ -634,6 +634,8 @@ GOACC_data_start (int flags_m, size_t mapnum, if (profiling_p) goacc_profiling_dispatch (&prof_info, &enter_data_event_info, &api_info); + handle_ftn_pointers (mapnum, hostaddrs, sizes, kinds); + /* Host fallback or 'do nothing'. */ if ((acc_dev->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM) || (flags & GOACC_FLAG_HOST_FALLBACK) diff --git a/libgomp/testsuite/libgomp.oacc-fortran/deviceptr-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/deviceptr-1.f90 new file mode 100644 index 000000000000..276a1727b2ee --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-fortran/deviceptr-1.f90 @@ -0,0 +1,197 @@ +! { dg-do run } + +! Test the deviceptr clause with various directives +! and in combination with other directives where +! the deviceptr variable is implied. + +subroutine subr1 (a, b) + implicit none + integer, parameter :: N = 8 + integer :: a(N) + integer :: b(N) + integer :: i = 0 + + !$acc data deviceptr (a) + + !$acc parallel copy (b) + do i = 1, N + a(i) = i * 2 + b(i) = a(i) + end do + !$acc end parallel + + !$acc end data + +end subroutine + +subroutine subr2 (a, b) + implicit none + integer, parameter :: N = 8 + integer :: a(N) + !$acc declare deviceptr (a) + integer :: b(N) + integer :: i = 0 + + !$acc parallel copy (b) + do i = 1, N + a(i) = i * 4 + b(i) = a(i) + end do + !$acc end parallel + +end subroutine + +subroutine subr3 (a, b) + implicit none + integer, parameter :: N = 8 + integer :: a(N) + !$acc declare deviceptr (a) + integer :: b(N) + integer :: i = 0 + + !$acc kernels copy (b) + do i = 1, N + a(i) = i * 8 + b(i) = a(i) + end do + !$acc end kernels + +end subroutine + +subroutine subr4 (a, b) + implicit none + integer, parameter :: N = 8 + integer :: a(N) + integer :: b(N) + integer :: i = 0 + + !$acc parallel deviceptr (a) copy (b) + do i = 1, N + a(i) = i * 16 + b(i) = a(i) + end do + !$acc end parallel + +end subroutine + +subroutine subr5 (a, b) + implicit none + integer, parameter :: N = 8 + integer :: a(N) + integer :: b(N) + integer :: i = 0 + + !$acc kernels deviceptr (a) copy (b) + do i = 1, N + a(i) = i * 32 + b(i) = a(i) + end do + !$acc end kernels + +end subroutine + +subroutine subr6 (a, b) + implicit none + integer, parameter :: N = 8 + integer :: a(N) + integer :: b(N) + integer :: i = 0 + + !$acc parallel deviceptr (a) copy (b) + do i = 1, N + b(i) = i + end do + !$acc end parallel + +end subroutine + +subroutine subr7 (a, b) + implicit none + integer, parameter :: N = 8 + integer :: a(N) + integer :: b(N) + integer :: i = 0 + + !$acc data deviceptr (a) + + !$acc parallel copy (b) + do i = 1, N + a(i) = i * 2 + b(i) = a(i) + end do + !$acc end parallel + + !$acc parallel copy (b) + do i = 1, N + a(i) = b(i) * 2 + b(i) = a(i) + end do + !$acc end parallel + + !$acc end data + +end subroutine + +program main + use iso_c_binding, only: c_ptr, c_f_pointer + implicit none + type (c_ptr) :: cp + integer, parameter :: N = 8 + integer, pointer :: fp(:) + integer :: i = 0 + integer :: b(N) + + interface + function acc_malloc (s) bind (C) + use iso_c_binding, only: c_ptr, c_size_t + integer (c_size_t), value :: s + type (c_ptr) :: acc_malloc + end function + end interface + + cp = acc_malloc (N * sizeof (fp(N))) + call c_f_pointer (cp, fp, [N]) + + call subr1 (fp, b) + + do i = 1, N + if (b(i) .ne. i * 2) call abort + end do + + call subr2 (fp, b) + + do i = 1, N + if (b(i) .ne. i * 4) call abort + end do + + call subr3 (fp, b) + + do i = 1, N + if (b(i) .ne. i * 8) call abort + end do + + call subr4 (fp, b) + + do i = 1, N + if (b(i) .ne. i * 16) call abort + end do + + call subr5 (fp, b) + + do i = 1, N + if (b(i) .ne. i * 32) call abort + end do + + call subr6 (fp, b) + + do i = 1, N + if (b(i) .ne. i) call abort + end do + + call subr7 (fp, b) + + do i = 1, N + if (b(i) .ne. i * 4) call abort + end do + +end program main