From 2d11254d63174644a80a869f031a3d929d7d4e4f Mon Sep 17 00:00:00 2001 From: Julian Brown Date: Tue, 12 Feb 2019 14:32:34 -0800 Subject: [PATCH] Add OpenACC Fortran support for deviceptr and variable in common blocks 2018-06-29 Cesar Philippidis James Norris gcc/fortran/ * openmp.c (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.c (gfc_omp_finish_clause): Don't create pointer data mappings for deviceptr clauses. (gfc_trans_omp_clauses): Likewise. gcc/ * gimplify.c (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/ * c-c++-common/goacc/deviceptr-4.c: Update. * gfortran.dg/goacc/common-block-1.f90: New test. * gfortran.dg/goacc/common-block-2.f90: New test. * gfortran.dg/goacc/loop-2.f95: Update. * gfortran.dg/goacc/loop-3-2.f95: Update. * gfortran.dg/goacc/loop-3.f95: Update. * gfortran.dg/goacc/loop-5.f95: Update. * gfortran.dg/goacc/pr72715.f90: New test. * gfortran.dg/goacc/sie.f95: Update. * gfortran.dg/goacc/tile-1.f90: Update. * gfortran.dg/gomp/pr77516.f90: Update. libgomp/ * oacc-parallel.c (GOACC_parallel_keyed): Handle Fortran deviceptr clause. (GOACC_data_start): Likewise. * testsuite/libgomp.oacc-fortran/common-block-1.f90: New test. * testsuite/libgomp.oacc-fortran/common-block-2.f90: New test. * testsuite/libgomp.oacc-fortran/common-block-3.f90: New test. * testsuite/libgomp.oacc-fortran/deviceptr-1.f90: New test. (cherry picked from openacc-gcc-9-branch commit ac6c90812344f4f4cfe4d2f5901c1a9d038a4000) --- gcc/ChangeLog.omp | 12 ++ gcc/fortran/ChangeLog.omp | 14 ++ gcc/fortran/openmp.c | 62 +++--- gcc/fortran/trans-openmp.c | 9 + gcc/gimplify.c | 35 +++- gcc/testsuite/ChangeLog.omp | 15 ++ .../c-c++-common/goacc/deviceptr-4.c | 2 +- .../gfortran.dg/goacc/common-block-1.f90 | 69 ++++++ .../gfortran.dg/goacc/common-block-2.f90 | 49 +++++ .../gfortran.dg/goacc/loop-2-kernels-tile.f95 | 4 +- .../goacc/loop-2-parallel-tile.f95 | 4 +- gcc/testsuite/gfortran.dg/goacc/loop-5.f95 | 12 -- gcc/testsuite/gfortran.dg/goacc/sie.f95 | 36 ++-- gcc/testsuite/gfortran.dg/goacc/tile-1.f90 | 16 +- gcc/testsuite/gfortran.dg/gomp/pr77516.f90 | 2 +- libgomp/ChangeLog.omp | 11 + libgomp/oacc-parallel.c | 2 + .../libgomp.oacc-fortran/common-block-1.f90 | 105 ++++++++++ .../libgomp.oacc-fortran/common-block-2.f90 | 150 +++++++++++++ .../libgomp.oacc-fortran/common-block-3.f90 | 137 ++++++++++++ .../libgomp.oacc-fortran/deviceptr-1.f90 | 197 ++++++++++++++++++ 21 files changed, 865 insertions(+), 78 deletions(-) create mode 100644 gcc/testsuite/gfortran.dg/goacc/common-block-1.f90 create mode 100644 gcc/testsuite/gfortran.dg/goacc/common-block-2.f90 create mode 100644 libgomp/testsuite/libgomp.oacc-fortran/common-block-1.f90 create mode 100644 libgomp/testsuite/libgomp.oacc-fortran/common-block-2.f90 create mode 100644 libgomp/testsuite/libgomp.oacc-fortran/common-block-3.f90 create mode 100644 libgomp/testsuite/libgomp.oacc-fortran/deviceptr-1.f90 diff --git a/gcc/ChangeLog.omp b/gcc/ChangeLog.omp index d50b791fa963..194aa754dca6 100644 --- a/gcc/ChangeLog.omp +++ b/gcc/ChangeLog.omp @@ -1,3 +1,15 @@ +2018-06-29 Cesar Philippidis + James Norris + + * gimplify.c (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. + 2018-10-02 Thomas Schwinge Cesar Philippidis diff --git a/gcc/fortran/ChangeLog.omp b/gcc/fortran/ChangeLog.omp index 933e32e68294..e36d18b3c889 100644 --- a/gcc/fortran/ChangeLog.omp +++ b/gcc/fortran/ChangeLog.omp @@ -1,3 +1,17 @@ +2018-06-29 Cesar Philippidis + James Norris + + * openmp.c (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.c (gfc_omp_finish_clause): Don't create pointer data + mappings for deviceptr clauses. + (gfc_trans_omp_clauses): Likewise. + 2018-10-02 Thomas Schwinge Cesar Philippidis diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c index bfceb4010e64..fa0a67399c04 100644 --- a/gcc/fortran/openmp.c +++ b/gcc/fortran/openmp.c @@ -920,10 +920,10 @@ omp_inv_mask::omp_inv_mask (const omp_mask &m) : omp_mask (m) static bool gfc_match_omp_map_clause (gfc_omp_namelist **list, gfc_omp_map_op map_op, - bool allow_derived = false) + bool common_blocks, bool allow_derived = false) { gfc_omp_namelist **head = NULL; - if (gfc_match_omp_variable_list ("", list, false, NULL, &head, true, + if (gfc_match_omp_variable_list ("", list, common_blocks, NULL, &head, true, allow_derived) == MATCH_YES) { @@ -1030,7 +1030,8 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask, if ((mask & OMP_CLAUSE_ATTACH) && gfc_match ("attach ( ") == MATCH_YES && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP], - OMP_MAP_ATTACH, allow_derived)) + OMP_MAP_ATTACH, false, + allow_derived)) continue; break; case 'c': @@ -1059,7 +1060,8 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask, if ((mask & OMP_CLAUSE_COPY) && gfc_match ("copy ( ") == MATCH_YES && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP], - OMP_MAP_TOFROM, allow_derived)) + OMP_MAP_TOFROM, openacc, + allow_derived)) continue; if (mask & OMP_CLAUSE_COPYIN) { @@ -1067,7 +1069,8 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask, { if (gfc_match ("copyin ( ") == MATCH_YES && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP], - OMP_MAP_TO, allow_derived)) + OMP_MAP_TO, true, + allow_derived)) continue; } else if (gfc_match_omp_variable_list ("copyin (", @@ -1078,7 +1081,7 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask, if ((mask & OMP_CLAUSE_COPYOUT) && gfc_match ("copyout ( ") == MATCH_YES && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP], - OMP_MAP_FROM, allow_derived)) + OMP_MAP_FROM, true, allow_derived)) continue; if ((mask & OMP_CLAUSE_COPYPRIVATE) && gfc_match_omp_variable_list ("copyprivate (", @@ -1088,7 +1091,7 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask, if ((mask & OMP_CLAUSE_CREATE) && gfc_match ("create ( ") == MATCH_YES && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP], - OMP_MAP_ALLOC, allow_derived)) + OMP_MAP_ALLOC, true, allow_derived)) continue; break; case 'd': @@ -1124,7 +1127,8 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask, if ((mask & OMP_CLAUSE_DELETE) && gfc_match ("delete ( ") == MATCH_YES && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP], - OMP_MAP_RELEASE, allow_derived)) + OMP_MAP_RELEASE, true, + allow_derived)) continue; if ((mask & OMP_CLAUSE_DEPEND) && gfc_match ("depend ( ") == MATCH_YES) @@ -1170,7 +1174,8 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask, if ((mask & OMP_CLAUSE_DETACH) && gfc_match ("detach ( ") == MATCH_YES && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP], - OMP_MAP_DETACH, allow_derived)) + OMP_MAP_DETACH, false, + allow_derived)) continue; if ((mask & OMP_CLAUSE_DEVICE) && !openacc @@ -1181,12 +1186,13 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask, && openacc && gfc_match ("device ( ") == MATCH_YES && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP], - OMP_MAP_FORCE_TO, allow_derived)) + OMP_MAP_FORCE_TO, false, + allow_derived)) continue; if ((mask & OMP_CLAUSE_DEVICEPTR) && gfc_match ("deviceptr ( ") == MATCH_YES && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP], - OMP_MAP_FORCE_DEVICEPTR, + OMP_MAP_FORCE_DEVICEPTR, false, allow_derived)) continue; if ((mask & OMP_CLAUSE_DEVICE_RESIDENT) @@ -1265,7 +1271,8 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask, if ((mask & OMP_CLAUSE_HOST_SELF) && gfc_match ("host ( ") == MATCH_YES && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP], - OMP_MAP_FORCE_FROM, allow_derived)) + OMP_MAP_FORCE_FROM, true, + allow_derived)) continue; break; case 'i': @@ -1544,48 +1551,50 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask, if ((mask & OMP_CLAUSE_COPY) && gfc_match ("pcopy ( ") == MATCH_YES && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP], - OMP_MAP_TOFROM, allow_derived)) + OMP_MAP_TOFROM, true, + allow_derived)) continue; if ((mask & OMP_CLAUSE_COPYIN) && gfc_match ("pcopyin ( ") == MATCH_YES && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP], - OMP_MAP_TO, allow_derived)) + OMP_MAP_TO, true, allow_derived)) continue; if ((mask & OMP_CLAUSE_COPYOUT) && gfc_match ("pcopyout ( ") == MATCH_YES && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP], - OMP_MAP_FROM, allow_derived)) + OMP_MAP_FROM, true, allow_derived)) continue; if ((mask & OMP_CLAUSE_CREATE) && gfc_match ("pcreate ( ") == MATCH_YES && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP], - OMP_MAP_ALLOC, allow_derived)) + OMP_MAP_ALLOC, true, allow_derived)) continue; if ((mask & OMP_CLAUSE_PRESENT) && gfc_match ("present ( ") == MATCH_YES && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP], - OMP_MAP_FORCE_PRESENT, + OMP_MAP_FORCE_PRESENT, false, allow_derived)) continue; if ((mask & OMP_CLAUSE_COPY) && gfc_match ("present_or_copy ( ") == MATCH_YES && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP], - OMP_MAP_TOFROM, allow_derived)) + OMP_MAP_TOFROM, true, + allow_derived)) continue; if ((mask & OMP_CLAUSE_COPYIN) && gfc_match ("present_or_copyin ( ") == MATCH_YES && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP], - OMP_MAP_TO, allow_derived)) + OMP_MAP_TO, true, allow_derived)) continue; if ((mask & OMP_CLAUSE_COPYOUT) && gfc_match ("present_or_copyout ( ") == MATCH_YES && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP], - OMP_MAP_FROM, allow_derived)) + OMP_MAP_FROM, true, allow_derived)) continue; if ((mask & OMP_CLAUSE_CREATE) && gfc_match ("present_or_create ( ") == MATCH_YES && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP], - OMP_MAP_ALLOC, allow_derived)) + OMP_MAP_ALLOC, true, allow_derived)) continue; if ((mask & OMP_CLAUSE_PRIORITY) && c->priority == NULL @@ -1803,7 +1812,8 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask, if ((mask & OMP_CLAUSE_HOST_SELF) && gfc_match ("self ( ") == MATCH_YES && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP], - OMP_MAP_FORCE_FROM, allow_derived)) + OMP_MAP_FORCE_FROM, true, + allow_derived)) continue; if ((mask & OMP_CLAUSE_SEQ) && !c->seq @@ -3821,8 +3831,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 (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 @@ -3877,10 +3887,6 @@ check_array_not_assumed (gfc_symbol *sym, locus loc, const char *name) if (sym->as && sym->as->type == AS_ASSUMED_RANK) gfc_error ("Assumed rank array %qs in %s clause at %L", sym->name, name, &loc); - if (sym->as && sym->as->type == AS_DEFERRED && sym->attr.pointer - && !sym->attr.contiguous) - gfc_error ("Noncontiguous deferred shape array %qs in %s clause at %L", - sym->name, name, &loc); } static void diff --git a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c index 6c30f0620ceb..3911945a2b36 100644 --- a/gcc/fortran/trans-openmp.c +++ b/gcc/fortran/trans-openmp.c @@ -1094,6 +1094,8 @@ gfc_omp_finish_clause (tree c, gimple_seq *pre_p) } tree c2 = NULL_TREE, c3 = NULL_TREE, c4 = NULL_TREE; + if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FORCE_DEVICEPTR) + return; if (POINTER_TYPE_P (TREE_TYPE (decl))) { if (!gfc_omp_privatize_by_reference (decl) @@ -2171,6 +2173,12 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, field = TREE_OPERAND (field, 1); if (POINTER_TYPE_P (TREE_TYPE (decl)) + && n->u.map_op == OMP_MAP_FORCE_DEVICEPTR) + { + OMP_CLAUSE_DECL (node) = decl; + goto finalize_map_clause; + } + else if (POINTER_TYPE_P (TREE_TYPE (decl)) && (gfc_omp_privatize_by_reference (decl) || GFC_DECL_GET_SCALAR_POINTER (field) || GFC_DECL_GET_SCALAR_ALLOCATABLE (field) @@ -2348,6 +2356,7 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, ptr2 = fold_convert (sizetype, ptr2); OMP_CLAUSE_SIZE (node3) = fold_build2 (MINUS_EXPR, sizetype, ptr, ptr2); + finalize_map_clause:; } else gcc_unreachable (); diff --git a/gcc/gimplify.c b/gcc/gimplify.c index ed9ceb633456..7927faf07a89 100644 --- a/gcc/gimplify.c +++ b/gcc/gimplify.c @@ -117,6 +117,9 @@ enum gimplify_omp_var_data fields. */ GOVD_MAP_HAS_ATTACHMENTS = 8388608, + /* Flag for OpenACC deviceptrs. */ + GOVD_DEVICEPTR = (1<<24), + GOVD_DATA_SHARE_CLASS = (GOVD_SHARED | GOVD_PRIVATE | GOVD_FIRSTPRIVATE | GOVD_LASTPRIVATE | GOVD_REDUCTION | GOVD_LINEAR | GOVD_LOCAL) @@ -7137,15 +7140,20 @@ oacc_default_clause (struct gimplify_omp_ctx *ctx, tree decl, unsigned flags) { const char *rkind; bool on_device = false; + bool is_private = false; bool declared = is_oacc_declared (decl); tree type = TREE_TYPE (decl); if (lang_hooks.decls.omp_privatize_by_reference (decl)) type = TREE_TYPE (type); + if (RECORD_OR_UNION_TYPE_P (type)) + is_private = lang_hooks.decls.omp_disregard_value_expr (decl, false); + if ((ctx->region_type & (ORT_ACC_PARALLEL | ORT_ACC_KERNELS)) != 0 && is_global_var (decl) - && device_resident_p (decl)) + && device_resident_p (decl) + && !is_private) { on_device = true; flags |= GOVD_MAP_TO_ONLY; @@ -7156,7 +7164,9 @@ oacc_default_clause (struct gimplify_omp_ctx *ctx, tree decl, unsigned flags) case ORT_ACC_KERNELS: rkind = "kernels"; - if (AGGREGATE_TYPE_P (type)) + if (is_private) + flags |= GOVD_MAP; + else if (AGGREGATE_TYPE_P (type)) { /* Aggregates default to 'present_or_copy', or 'present'. */ if (ctx->default_kind != OMP_CLAUSE_DEFAULT_PRESENT) @@ -7173,7 +7183,9 @@ oacc_default_clause (struct gimplify_omp_ctx *ctx, tree decl, unsigned flags) case ORT_ACC_PARALLEL: rkind = "parallel"; - if (on_device || declared) + if (is_private) + flags |= GOVD_FIRSTPRIVATE; + else if (on_device || declared) flags |= GOVD_MAP; else if (AGGREGATE_TYPE_P (type)) { @@ -7239,7 +7251,8 @@ omp_notice_variable (struct gimplify_omp_ctx *ctx, tree decl, bool in_code) { tree value = get_base_address (DECL_VALUE_EXPR (decl)); - if (value && DECL_P (value) && DECL_THREAD_LOCAL_P (value)) + if (!(ctx->region_type & ORT_ACC) + && value && DECL_P (value) && DECL_THREAD_LOCAL_P (value)) return omp_notice_threadprivate_variable (ctx, decl, value); } @@ -7271,7 +7284,8 @@ omp_notice_variable (struct gimplify_omp_ctx *ctx, tree decl, bool in_code) n = splay_tree_lookup (ctx->variables, (splay_tree_key)decl); if ((ctx->region_type & ORT_TARGET) != 0) { - ret = lang_hooks.decls.omp_disregard_value_expr (decl, true); + shared = !(ctx->region_type & ORT_ACC); + ret = lang_hooks.decls.omp_disregard_value_expr (decl, shared); if (n == NULL) { unsigned nflags = flags; @@ -7344,6 +7358,7 @@ omp_notice_variable (struct gimplify_omp_ctx *ctx, tree decl, bool in_code) error ("variable %qE declared in enclosing " "% 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; @@ -7439,6 +7454,8 @@ omp_notice_variable (struct gimplify_omp_ctx *ctx, tree decl, bool in_code) } shared = ((flags | n->value) & GOVD_SHARED) != 0; + if (ctx->region_type & ORT_ACC) + shared = false; ret = lang_hooks.decls.omp_disregard_value_expr (decl, shared); /* If nothing changed, there's nothing left to do. */ @@ -9008,6 +9025,8 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_TO || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_TOFROM) flags |= GOVD_MAP_ALWAYS_TO; + else if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FORCE_DEVICEPTR) + flags |= GOVD_DEVICEPTR; goto do_add; case OMP_CLAUSE_DEPEND: @@ -9743,7 +9762,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; @@ -9766,6 +9786,9 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n, void *data) case GOVD_MAP_FORCE_PRESENT: kind = GOMP_MAP_FORCE_PRESENT; break; + case GOVD_DEVICEPTR: + kind = GOMP_MAP_FORCE_DEVICEPTR; + break; default: gcc_unreachable (); } diff --git a/gcc/testsuite/ChangeLog.omp b/gcc/testsuite/ChangeLog.omp index adad65edb20f..e379c43cffe5 100644 --- a/gcc/testsuite/ChangeLog.omp +++ b/gcc/testsuite/ChangeLog.omp @@ -1,3 +1,18 @@ +2018-06-29 Cesar Philippidis + James Norris + + * c-c++-common/goacc/deviceptr-4.c: Update. + * gfortran.dg/goacc/common-block-1.f90: New test. + * gfortran.dg/goacc/common-block-2.f90: New test. + * gfortran.dg/goacc/loop-2.f95: Update. + * gfortran.dg/goacc/loop-3-2.f95: Update. + * gfortran.dg/goacc/loop-3.f95: Update. + * gfortran.dg/goacc/loop-5.f95: Update. + * gfortran.dg/goacc/pr72715.f90: New test. + * gfortran.dg/goacc/sie.f95: Update. + * gfortran.dg/goacc/tile-1.f90: Update. + * gfortran.dg/gomp/pr77516.f90: Update. + 2018-10-02 Thomas Schwinge Cesar Philippidis 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/common-block-1.f90 b/gcc/testsuite/gfortran.dg/goacc/common-block-1.f90 new file mode 100644 index 000000000000..c9de125a2f17 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/goacc/common-block-1.f90 @@ -0,0 +1,69 @@ +! Test data clauses involving common blocks and common block data. +! Specifically, validate early matching errors. + +subroutine subtest + implicit none + integer, parameter :: n = 10 + integer a(n), b(n), c, d(n), e + real*4 x(n), y(n), z, w(n), v + common /blockA/ a, c, x + common /blockB/ b, y, z + !$acc declare link(/blockA/, /blockB/, e, v) +end subroutine subtest + +program test + implicit none + integer, parameter :: n = 10 + integer a(n), b(n), c, d(n), e + real*4 x(n), y(n), z, w(n), v + common /blockA/ a, c, x + common /blockB/ b, y, z + !$acc declare link(/blockA/, /blockB/, e, v) + + !$acc data copy(/blockA/, /blockB/, e, v) + !$acc end data + + !$acc data copyin(/blockA/, /blockB/, e, v) + !$acc end data + + !$acc data copyout(/blockA/, /blockB/, e, v) + !$acc end data + + !$acc data create(/blockA/, /blockB/, e, v) + !$acc end data + + !$acc data copyout(/blockA/, /blockB/, e, v) + !$acc end data + + !$acc data pcopy(/blockA/, /blockB/, e, v) + !$acc end data + + !$acc data pcopyin(/blockA/, /blockB/, e, v) + !$acc end data + + !$acc data pcopyout(/blockA/, /blockB/, e, v) + !$acc end data + + !$acc data pcreate(/blockA/, /blockB/, e, v) + !$acc end data + + !$acc data pcopyout(/blockA/, /blockB/, e, v) + !$acc end data + + !$acc data present(/blockA/, /blockB/, e, v) ! { dg-error "" } + !$acc end data ! { dg-error "Unexpected ..ACC END DATA statement" } + + !$acc parallel private(/blockA/, /blockB/, e, v) + !$acc end parallel + + !$acc parallel firstprivate(/blockA/, /blockB/, e, v) + !$acc end parallel + + !$acc exit data delete(/blockA/, /blockB/, e, v) + + !$acc data deviceptr(/blockA/, /blockB/, e, v) ! { dg-error "Syntax error" } + !$acc end data ! { dg-error "Unexpected ..ACC END DATA statement" } + + !$acc data deviceptr(/blockA/, /blockB/, e, v) ! { dg-error "Syntax error" } + !$acc end data ! { dg-error "Unexpected ..ACC END DATA statement" } +end program test diff --git a/gcc/testsuite/gfortran.dg/goacc/common-block-2.f90 b/gcc/testsuite/gfortran.dg/goacc/common-block-2.f90 new file mode 100644 index 000000000000..b83638918a3d --- /dev/null +++ b/gcc/testsuite/gfortran.dg/goacc/common-block-2.f90 @@ -0,0 +1,49 @@ +! Test data clauses involving common blocks and common block data. +! Specifically, resolver errors such as duplicate data clauses. + +program test + implicit none + integer, parameter :: n = 10 + integer a(n), b(n), c, d(n), e + real*4 x(n), y(n), z, w(n), v + common /blockA/ a, c, x + common /blockB/ b, y, z + + !$acc data copy(/blockA/, /blockB/, e, v, a) ! { dg-error "Symbol .a. present on multiple clauses" } + !$acc end data + + !$acc data copyin(/blockA/, /blockB/, e, v, a) ! { dg-error "Symbol .a. present on multiple clauses" } + !$acc end data + + !$acc data copyout(/blockA/, /blockB/, e, v, a) ! { dg-error "Symbol .a. present on multiple clauses" } + !$acc end data + + !$acc data create(/blockA/, /blockB/, e, v, a) ! { dg-error "Symbol .a. present on multiple clauses" } + !$acc end data + + !$acc data copyout(/blockA/, /blockB/, e, v, a) ! { dg-error "Symbol .a. present on multiple clauses" } + !$acc end data + + !$acc data pcopy(/blockA/, /blockB/, e, v, a) ! { dg-error "Symbol .a. present on multiple clauses" } + !$acc end data + + !$acc data pcopyin(/blockA/, /blockB/, e, v, a) ! { dg-error "Symbol .a. present on multiple clauses" } + !$acc end data + + !$acc data pcopyout(/blockA/, /blockB/, e, v, a) ! { dg-error "Symbol .a. present on multiple clauses" } + !$acc end data + + !$acc data pcreate(/blockA/, /blockB/, e, v, a) ! { dg-error "Symbol .a. present on multiple clauses" } + !$acc end data + + !$acc data pcopyout(/blockA/, /blockB/, e, v, a) ! { dg-error "Symbol .a. present on multiple clauses" } + !$acc end data + + !$acc parallel private(/blockA/, /blockB/, e, v, a) ! { dg-error "Symbol .a. present on multiple clauses" } + !$acc end parallel + + !$acc parallel firstprivate(/blockA/, /blockB/, e, v, a) ! { dg-error "Symbol .a. present on multiple clauses" } + !$acc end parallel + + !$acc exit data delete(/blockA/, /blockB/, e, v, a) ! { dg-error "Symbol .a. present on multiple clauses" } +end program test 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 96e0ccbd2ac7..937cd73f6ee8 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 3a4db5dc1afc..e942011c6009 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-5.f95 b/gcc/testsuite/gfortran.dg/goacc/loop-5.f95 index d059cf7f3777..fe137d515ee6 100644 --- a/gcc/testsuite/gfortran.dg/goacc/loop-5.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/loop-5.f95 @@ -93,9 +93,6 @@ program test DO j = 1,10 ENDDO ENDDO - !$acc loop tile(-1) ! { dg-warning "must be positive" } - do i = 1,10 - enddo !$acc loop vector tile(*) DO i = 1,10 ENDDO @@ -129,9 +126,6 @@ program test DO j = 1,10 ENDDO ENDDO - !$acc loop tile(-1) ! { dg-warning "must be positive" } - do i = 1,10 - enddo !$acc loop vector tile(*) DO i = 1,10 ENDDO @@ -242,9 +236,6 @@ program test DO j = 1,10 ENDDO ENDDO - !$acc kernels loop tile(-1) ! { dg-warning "must be positive" } - do i = 1,10 - enddo !$acc kernels loop vector tile(*) DO i = 1,10 ENDDO @@ -333,9 +324,6 @@ program test DO j = 1,10 ENDDO ENDDO - !$acc parallel loop tile(-1) ! { dg-warning "must be positive" } - do i = 1,10 - enddo !$acc parallel loop vector tile(*) DO i = 1,10 ENDDO diff --git a/gcc/testsuite/gfortran.dg/goacc/sie.f95 b/gcc/testsuite/gfortran.dg/goacc/sie.f95 index abfe28bc5337..3abf2c8016f5 100644 --- a/gcc/testsuite/gfortran.dg/goacc/sie.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/sie.f95 @@ -78,10 +78,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" } @@ -106,10 +106,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" } @@ -135,10 +135,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 character in name" } @@ -163,10 +163,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 character in name" } @@ -192,10 +192,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 character in name" } @@ -220,10 +220,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 character in name" } @@ -250,10 +250,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" } @@ -282,10 +282,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" } @@ -314,10 +314,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 3dbabda0342e..17fd32cd2841 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/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/ChangeLog.omp b/libgomp/ChangeLog.omp index 2da291cca004..9f5e057350c6 100644 --- a/libgomp/ChangeLog.omp +++ b/libgomp/ChangeLog.omp @@ -1,3 +1,14 @@ +2018-06-29 Cesar Philippidis + James Norris + + * oacc-parallel.c (GOACC_parallel_keyed): Handle Fortran deviceptr + clause. + (GOACC_data_start): Likewise. + * testsuite/libgomp.oacc-fortran/common-block-1.f90: New test. + * testsuite/libgomp.oacc-fortran/common-block-2.f90: New test. + * testsuite/libgomp.oacc-fortran/common-block-3.f90: New test. + * testsuite/libgomp.oacc-fortran/deviceptr-1.f90: New test. + 2019-02-12 Julian Brown * oacc-cuda.c (acc_set_cuda_stream): Return 0 on error/invalid diff --git a/libgomp/oacc-parallel.c b/libgomp/oacc-parallel.c index f588d7e42b5a..8407e3b34bc6 100644 --- a/libgomp/oacc-parallel.c +++ b/libgomp/oacc-parallel.c @@ -327,6 +327,8 @@ GOACC_data_start (int flags_m, size_t mapnum, struct goacc_thread *thr = goacc_thread (); struct gomp_device_descr *acc_dev = thr->dev; + 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/common-block-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/common-block-1.f90 new file mode 100644 index 000000000000..9f402973d3df --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-fortran/common-block-1.f90 @@ -0,0 +1,105 @@ +! Test data located inside common blocks. This test does not execrise +! ACC DECLARE. + +module const + integer, parameter :: n = 100 +end module const + +subroutine check + use const + + implicit none + integer i, x(n), y + common /BLOCK/ x, y + + do i = 1, n + if (x(i) .ne. y) call abort + end do +end subroutine check + +module m + use const + integer a(n), b + common /BLOCK/ a, b + +contains + subroutine mod_implicit_incr + implicit none + integer i + + !$acc parallel loop + do i = 1, n + a(i) = b + end do + !$acc end parallel loop + + call check + end subroutine mod_implicit_incr + + subroutine mod_explicit_incr + implicit none + integer i + + !$acc parallel loop copy(a(1:n)) copyin(b) + do i = 1, n + a(i) = b + end do + !$acc end parallel loop + + call check + end subroutine mod_explicit_incr +end module m + +subroutine sub_implicit_incr + use const + + implicit none + integer i, x(n), y + common /BLOCK/ x, y + + !$acc parallel loop + do i = 1, n + x(i) = y + end do + !$acc end parallel loop + + call check +end subroutine sub_implicit_incr + +subroutine sub_explicit_incr + use const + + implicit none + integer i, x(n), y + common /BLOCK/ x, y + + !$acc parallel loop copy(x(1:n)) copyin(y) + do i = 1, n + x(i) = y + end do + !$acc end parallel loop + + call check +end subroutine sub_explicit_incr + +program main + use m + + implicit none + + a(:) = -1 + b = 5 + call mod_implicit_incr + + a(:) = -2 + b = 6 + call mod_explicit_incr + + a(:) = -3 + b = 7 + call sub_implicit_incr + + a(:) = -4 + b = 8 + call sub_explicit_incr +end program main diff --git a/libgomp/testsuite/libgomp.oacc-fortran/common-block-2.f90 b/libgomp/testsuite/libgomp.oacc-fortran/common-block-2.f90 new file mode 100644 index 000000000000..bf17fc586b90 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-fortran/common-block-2.f90 @@ -0,0 +1,150 @@ +! Test data located inside common blocks. This test does not execrise +! ACC DECLARE. All data clauses are explicit. + +module consts + integer, parameter :: n = 100 +end module consts + +subroutine validate + use consts + + implicit none + integer i, j + real*4 x(n), y(n), z + common /BLOCK/ x, y, z, j + + do i = 1, n + if (abs(x(i) - i - z) .ge. 0.0001) call abort + end do +end subroutine validate + +subroutine incr + use consts + + implicit none + integer i, j + real*4 x(n), y(n), z + common /BLOCK/ x, y, z, j + + !$acc parallel loop pcopy(/BLOCK/) + do i = 1, n + x(i) = x(i) + z + end do + !$acc end parallel loop +end subroutine incr + +program main + use consts + + implicit none + integer i, j + real*4 a(n), b(n), c + common /BLOCK/ a, b, c, j + + ! Test copyout, pcopy, device + + !$acc data copyout(a, c) + + c = 1.0 + + !$acc update device(c) + + !$acc parallel loop pcopy(a) + do i = 1, n + a(i) = i + end do + !$acc end parallel loop + + call incr + call incr + call incr + !$acc end data + + c = 3.0 + call validate + + ! Test pcopy without copyout + + c = 2.0 + call incr + c = 5.0 + call validate + + ! Test create, delete, host, copyout, copyin + + !$acc enter data create(b) + + !$acc parallel loop pcopy(b) + do i = 1, n + b(i) = i + end do + !$acc end parallel loop + + !$acc update host (b) + + !$acc parallel loop pcopy(b) copyout(a) copyin(c) + do i = 1, n + a(i) = b(i) + c + end do + !$acc end parallel loop + + !$acc exit data delete(b) + + call validate + + a(:) = b(:) + c = 0.0 + call validate + + ! Test copy + + c = 1.0 + !$acc parallel loop copy(/BLOCK/) + do i = 1, n + a(i) = b(i) + c + end do + !$acc end parallel loop + + call validate + + ! Test pcopyin, pcopyout FIXME + + c = 2.0 + !$acc data copyin(b, c) copyout(a) + + !$acc parallel loop pcopyin(b, c) pcopyout(a) + do i = 1, n + a(i) = b(i) + c + end do + !$acc end parallel loop + + !$acc end data + + call validate + + ! Test reduction, private + + j = 0 + + !$acc parallel private(i) copy(j) + !$acc loop reduction(+:j) + do i = 1, n + j = j + 1 + end do + !$acc end parallel + + if (j .ne. n) call abort + + ! Test firstprivate, copy + + a(:) = 0 + c = j + + !$acc parallel loop firstprivate(c) copyout(a) + do i = 1, n + a(i) = i + c + end do + !$acc end parallel loop + + call validate +end program main diff --git a/libgomp/testsuite/libgomp.oacc-fortran/common-block-3.f90 b/libgomp/testsuite/libgomp.oacc-fortran/common-block-3.f90 new file mode 100644 index 000000000000..134e2d1cf29b --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-fortran/common-block-3.f90 @@ -0,0 +1,137 @@ +! Test data located inside common blocks. This test does not execrise +! ACC DECLARE. Most of the data clauses are implicit. + +module consts + integer, parameter :: n = 100 +end module consts + +subroutine validate + use consts + + implicit none + integer i, j + real*4 x(n), y(n), z + common /BLOCK/ x, y, z, j + + do i = 1, n + if (abs(x(i) - i - z) .ge. 0.0001) call abort + end do +end subroutine validate + +subroutine incr_parallel + use consts + + implicit none + integer i, j + real*4 x(n), y(n), z + common /BLOCK/ x, y, z, j + + !$acc parallel loop + do i = 1, n + x(i) = x(i) + z + end do + !$acc end parallel loop +end subroutine incr_parallel + +subroutine incr_kernels + use consts + + implicit none + integer i, j + real*4 x(n), y(n), z + common /BLOCK/ x, y, z, j + + !$acc kernels + do i = 1, n + x(i) = x(i) + z + end do + !$acc end kernels +end subroutine incr_kernels + +program main + use consts + + implicit none + integer i, j + real*4 a(n), b(n), c + common /BLOCK/ a, b, c, j + + !$acc data copyout(a, c) + + c = 1.0 + + !$acc update device(c) + + !$acc parallel loop + do i = 1, n + a(i) = i + end do + !$acc end parallel loop + + call incr_parallel + call incr_parallel + call incr_parallel + !$acc end data + + c = 3.0 + call validate + + ! Test pcopy without copyout + + c = 2.0 + call incr_kernels + c = 5.0 + call validate + + !$acc kernels + do i = 1, n + b(i) = i + end do + !$acc end kernels + + !$acc parallel loop + do i = 1, n + a(i) = b(i) + c + end do + !$acc end parallel loop + + call validate + + a(:) = b(:) + c = 0.0 + call validate + + ! Test copy + + c = 1.0 + !$acc parallel loop + do i = 1, n + a(i) = b(i) + c + end do + !$acc end parallel loop + + call validate + + c = 2.0 + !$acc data copyin(b, c) copyout(a) + + !$acc kernels + do i = 1, n + a(i) = b(i) + c + end do + !$acc end kernels + + !$acc end data + + call validate + + j = 0 + + !$acc parallel loop reduction(+:j) + do i = 1, n + j = j + 1 + end do + !$acc end parallel loop + + if (j .ne. n) call abort +end program main 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 -- 2.47.2