From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: (qmail 30098 invoked by alias); 29 Jun 2018 19:05:12 -0000 Mailing-List: contact fortran-help@gcc.gnu.org; run by ezmlm Precedence: bulk List-Id: List-Subscribe: List-Post: List-Help: , Sender: fortran-owner@gcc.gnu.org Received: (qmail 30075 invoked by uid 89); 29 Jun 2018 19:05:12 -0000 Authentication-Results: sourceware.org; auth=none X-Spam-SWARE-Status: No, score=-25.3 required=5.0 tests=AWL,BAYES_00,GIT_PATCH_0,GIT_PATCH_1,GIT_PATCH_2,GIT_PATCH_3,RCVD_IN_DNSWL_NONE,SPF_PASS,URIBL_RED autolearn=ham version=3.3.2 spammy=mandate, gfortran's, gfortrans, rank X-HELO: relay1.mentorg.com Received: from relay1.mentorg.com (HELO relay1.mentorg.com) (192.94.38.131) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with ESMTP; Fri, 29 Jun 2018 19:05:03 +0000 Received: from svr-orw-mbx-01.mgc.mentorg.com ([147.34.90.201]) by relay1.mentorg.com with esmtps (TLSv1.2:ECDHE-RSA-AES256-SHA384:256) id 1fYyhZ-0001Yv-Sn from Cesar_Philippidis@mentor.com ; Fri, 29 Jun 2018 12:05:01 -0700 Received: from [127.0.0.1] (147.34.91.1) by svr-orw-mbx-01.mgc.mentorg.com (147.34.90.201) with Microsoft SMTP Server (TLS) id 15.0.1320.4; Fri, 29 Jun 2018 12:04:58 -0700 From: Cesar Philippidis Subject: [patch] Add OpenACC Fortran support for deviceptr and variable in common blocks To: Fortran List , "gcc-patches@gcc.gnu.org" , Jakub Jelinek Message-ID: Date: Sat, 30 Jun 2018 14:37:00 -0000 User-Agent: Mozilla/5.0 (X11; Linux x86_64; rv:52.0) Gecko/20100101 Thunderbird/52.8.0 MIME-Version: 1.0 Content-Type: multipart/mixed; boundary="------------196CD9E12987FD434A40FEB8" X-IsSubscribed: yes X-SW-Source: 2018-06/txt/msg00222.txt.bz2 --------------196CD9E12987FD434A40FEB8 Content-Type: text/plain; charset="utf-8" Content-Transfer-Encoding: 7bit Content-length: 1583 The attached patch adds support Fortran support for OpenACC deviceptr and the use of common block variables in data clauses (both implicit and explicit). This patch also relaxes the Fortran parser to not error certain types of integral expressions and assumed-sized arrays. With respect to those errors, I removed them because a lot of working applications do not explicitly use type attributes (like contiguous). Perhaps it would be better to reduce them to a warning. Any thoughts on that? My argument for their removal is that, while the standard states that, say, arrays must be contiguous or bad things will happen, it does not necessary mandate that the compiler enforces it. I.e., the intent is to set the user's expectation that things will go bad if garbage input is fed to the accelerator. If necessary, I can push back on the OpenACC standards committee on these issue, but don't expect a quick resolution. In hindsight, I probably should have kept the error relaxation patches separate. This patch includes the following patches from og8: * (dd8b75a) [OpenACC] Update deviceptr handling * (634727d) [OpenACC] Handle Fortran deviceptr clause * (d50862a) [Fortran] Remove pointer check in check_array_not_assumed * (0793cef) [OpenACC] add support for fortran common blocks * (bdc1acc) [Fortran] update gfortran's tile clause error handling * (5dc4968) Fix PR72715 "ICE in gfc_trans_omp_do, at fortran/trans-openmp.c:3164" Is this patch OK for trunk? It bootstrapped / regression tested cleanly for x86_64 with nvptx offloading. Thanks, Cesar --------------196CD9E12987FD434A40FEB8 Content-Type: text/x-patch; name="0001-deviceptr-patch-56.patch" Content-Transfer-Encoding: 7bit Content-Disposition: attachment; filename="0001-deviceptr-patch-56.patch" Content-length: 45517 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. >From 09c1aa87d9a7db2e08384bb47c80b4a61d218a99 Mon Sep 17 00:00:00 2001 From: Cesar Philippidis Date: Mon, 25 Jun 2018 13:10:13 -0700 Subject: [PATCH] fortran deviceptr dd8b75 [OpenACC] Update deviceptr handling 634727 [OpenACC] Handle Fortran deviceptr clause 0793ce [OpenACC] add support for fortran common blocks bdc1ac [Fortran] update gfortran's tile clause error handling d50862 [Fortran] Remove pointer check in check_array_not_assumed 5dc496 Fix PR72715 "ICE in gfc_trans_omp_do, at fortran/trans-openmp.c:3164" --- gcc/fortran/openmp.c | 57 ++--- gcc/fortran/trans-openmp.c | 9 + gcc/gimplify.c | 35 +++- .../c-c++-common/goacc/deviceptr-4.c | 2 +- .../gfortran.dg/goacc/common-block-1.f90 | 69 ++++++ .../gfortran.dg/goacc/common-block-2.f90 | 49 +++++ gcc/testsuite/gfortran.dg/goacc/loop-2.f95 | 8 +- gcc/testsuite/gfortran.dg/goacc/loop-3-2.f95 | 4 +- gcc/testsuite/gfortran.dg/goacc/loop-3.f95 | 4 +- gcc/testsuite/gfortran.dg/goacc/loop-5.f95 | 12 -- gcc/testsuite/gfortran.dg/goacc/pr72715.f90 | 6 + 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/oacc-parallel.c | 11 +- .../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 ++++++++++++++++++ 19 files changed, 826 insertions(+), 83 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 gcc/testsuite/gfortran.dg/goacc/pr72715.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/fortran/openmp.c b/gcc/fortran/openmp.c index 94a7f7eaa50..aec2b26deef 100644 --- a/gcc/fortran/openmp.c +++ b/gcc/fortran/openmp.c @@ -914,10 +914,11 @@ omp_inv_mask::omp_inv_mask (const omp_mask &m) : omp_mask (m) mapping. */ static bool -gfc_match_omp_map_clause (gfc_omp_namelist **list, gfc_omp_map_op map_op) +gfc_match_omp_map_clause (gfc_omp_namelist **list, gfc_omp_map_op map_op, + bool common_blocks) { 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) == MATCH_YES) { gfc_omp_namelist *n; @@ -1039,7 +1040,7 @@ 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)) + OMP_MAP_TOFROM, openacc)) continue; if (mask & OMP_CLAUSE_COPYIN) { @@ -1047,7 +1048,7 @@ 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)) + OMP_MAP_TO, true)) continue; } else if (gfc_match_omp_variable_list ("copyin (", @@ -1058,7 +1059,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)) + OMP_MAP_FROM, true)) continue; if ((mask & OMP_CLAUSE_COPYPRIVATE) && gfc_match_omp_variable_list ("copyprivate (", @@ -1068,7 +1069,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)) + OMP_MAP_ALLOC, true)) continue; break; case 'd': @@ -1104,7 +1105,7 @@ 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)) + OMP_MAP_RELEASE, true)) continue; if ((mask & OMP_CLAUSE_DEPEND) && gfc_match ("depend ( ") == MATCH_YES) @@ -1156,12 +1157,12 @@ 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)) + OMP_MAP_FORCE_TO, false)) 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)) continue; if ((mask & OMP_CLAUSE_DEVICE_RESIDENT) && gfc_match_omp_variable_list @@ -1239,7 +1240,7 @@ 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)) + OMP_MAP_FORCE_FROM, true)) continue; break; case 'i': @@ -1511,47 +1512,47 @@ 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)) + OMP_MAP_TOFROM, true)) continue; if ((mask & OMP_CLAUSE_COPYIN) && gfc_match ("pcopyin ( ") == MATCH_YES && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP], - OMP_MAP_TO)) + OMP_MAP_TO, true)) continue; if ((mask & OMP_CLAUSE_COPYOUT) && gfc_match ("pcopyout ( ") == MATCH_YES && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP], - OMP_MAP_FROM)) + OMP_MAP_FROM, true)) continue; if ((mask & OMP_CLAUSE_CREATE) && gfc_match ("pcreate ( ") == MATCH_YES && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP], - OMP_MAP_ALLOC)) + OMP_MAP_ALLOC, true)) 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)) 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)) + OMP_MAP_TOFROM, true)) 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)) + OMP_MAP_TO, true)) 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)) + OMP_MAP_FROM, true)) 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)) + OMP_MAP_ALLOC, true)) continue; if ((mask & OMP_CLAUSE_PRIORITY) && c->priority == NULL @@ -1774,7 +1775,7 @@ 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)) + OMP_MAP_FORCE_FROM, true)) continue; if ((mask & OMP_CLAUSE_SEQ) && !c->seq @@ -3718,8 +3719,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 @@ -3777,10 +3778,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 @@ -5751,7 +5748,13 @@ resolve_oacc_nested_loops (gfc_code *code, gfc_code* do_code, int collapse, "at %L", &do_code->loc); break; } - gcc_assert (do_code->op == EXEC_DO || do_code->op == EXEC_DO_CONCURRENT); + if (do_code->op == EXEC_DO_CONCURRENT) + { + gfc_error ("!$ACC LOOP cannot be a DO CONCURRENT loop at %L", + &do_code->loc); + break; + } + gcc_assert (do_code->op == EXEC_DO); if (do_code->ext.iterator->var->ts.type != BT_INTEGER) gfc_error ("!$ACC LOOP iteration variable must be of type integer at %L", &do_code->loc); diff --git a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c index f038f4c5bf8..ca31c88569d 100644 --- a/gcc/fortran/trans-openmp.c +++ b/gcc/fortran/trans-openmp.c @@ -1060,6 +1060,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) @@ -2111,6 +2113,12 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, if (n->expr == NULL || n->expr->ref->u.ar.type == AR_FULL) { 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 (decl) || GFC_DECL_GET_SCALAR_ALLOCATABLE (decl) @@ -2282,6 +2290,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:; } switch (n->u.map_op) { diff --git a/gcc/gimplify.c b/gcc/gimplify.c index 48ac92e2b16..346b32031ad 100644 --- a/gcc/gimplify.c +++ b/gcc/gimplify.c @@ -105,6 +105,9 @@ enum gimplify_omp_var_data /* Flag for GOVD_MAP: must be present already. */ GOVD_MAP_FORCE_PRESENT = 524288, + /* Flag for OpenACC deviceptrs. */ + GOVD_DEVICEPTR = (1<<21), + GOVD_DATA_SHARE_CLASS = (GOVD_SHARED | GOVD_PRIVATE | GOVD_FIRSTPRIVATE | GOVD_LASTPRIVATE | GOVD_REDUCTION | GOVD_LINEAR | GOVD_LOCAL) @@ -7039,15 +7042,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; @@ -7058,7 +7066,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) @@ -7075,7 +7085,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)) { @@ -7141,7 +7153,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); } @@ -7173,7 +7186,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; @@ -7232,6 +7246,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; @@ -7326,6 +7341,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. */ @@ -8213,6 +8230,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: @@ -8828,7 +8847,8 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n, void *data) /* Not all combinations of these GOVD_MAP flags are actually valid. */ switch (flags & (GOVD_MAP_TO_ONLY | GOVD_MAP_FORCE - | GOVD_MAP_FORCE_PRESENT)) + | GOVD_MAP_FORCE_PRESENT + | GOVD_DEVICEPTR)) { case 0: kind = GOMP_MAP_TOFROM; @@ -8845,6 +8865,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/c-c++-common/goacc/deviceptr-4.c b/gcc/testsuite/c-c++-common/goacc/deviceptr-4.c index db1b91633a6..79a51620db9 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 00000000000..c9de125a2f1 --- /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 00000000000..b83638918a3 --- /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.f95 b/gcc/testsuite/gfortran.dg/goacc/loop-2.f95 index 0c902b22410..d4c62732331 100644 --- a/gcc/testsuite/gfortran.dg/goacc/loop-2.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/loop-2.f95 @@ -143,7 +143,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" } @@ -307,7 +307,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" } @@ -460,7 +460,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" } @@ -612,7 +612,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-3-2.f95 b/gcc/testsuite/gfortran.dg/goacc/loop-3-2.f95 index 9be74a85919..c091084a4f5 100644 --- a/gcc/testsuite/gfortran.dg/goacc/loop-3-2.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/loop-3-2.f95 @@ -27,9 +27,9 @@ subroutine test1 !$acc end parallel !$acc end loop ! { dg-error "Unexpected" } - ! OpenACC supports Fortran 2008 do concurrent statement + ! OpenACC does not support Fortran 2008 do concurrent statement !$acc loop - do concurrent (i = 1:5) + do concurrent (i = 1:5) ! { dg-error "ACC LOOP cannot be a DO CONCURRENT loop" } end do !$acc loop diff --git a/gcc/testsuite/gfortran.dg/goacc/loop-3.f95 b/gcc/testsuite/gfortran.dg/goacc/loop-3.f95 index 30930f404f3..ed3e8d50a76 100644 --- a/gcc/testsuite/gfortran.dg/goacc/loop-3.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/loop-3.f95 @@ -24,9 +24,9 @@ subroutine test1 !$acc end parallel !$acc end loop ! { dg-error "Unexpected" } - ! OpenACC supports Fortran 2008 do concurrent statement + ! OpenACC does not support Fortran 2008 do concurrent statement !$acc loop - do concurrent (i = 1:5) + do concurrent (i = 1:5) ! { dg-error "ACC LOOP cannot be a DO CONCURRENT loop" } end do !$acc loop diff --git a/gcc/testsuite/gfortran.dg/goacc/loop-5.f95 b/gcc/testsuite/gfortran.dg/goacc/loop-5.f95 index d059cf7f377..fe137d515ee 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/pr72715.f90 b/gcc/testsuite/gfortran.dg/goacc/pr72715.f90 new file mode 100644 index 00000000000..68580f9f7de --- /dev/null +++ b/gcc/testsuite/gfortran.dg/goacc/pr72715.f90 @@ -0,0 +1,6 @@ +program p + integer :: i + !$acc loop + do concurrent (i=1:3) ! { dg-error "ACC LOOP cannot be a DO CONCURRENT loop" } + end do +end program p diff --git a/gcc/testsuite/gfortran.dg/goacc/sie.f95 b/gcc/testsuite/gfortran.dg/goacc/sie.f95 index abfe28bc533..3abf2c8016f 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 3dbabda0342..17fd32cd284 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 9c0a95b9f79..3ac3f5562d0 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 b80ace58590..f9cfc4c3a16 100644 --- a/libgomp/oacc-parallel.c +++ b/libgomp/oacc-parallel.c @@ -231,8 +231,13 @@ GOACC_parallel_keyed (int device, void (*fn) (void *), devaddrs = gomp_alloca (sizeof (void *) * mapnum); for (i = 0; i < mapnum; i++) - devaddrs[i] = (void *) (tgt->list[i].key->tgt->tgt_start - + tgt->list[i].key->tgt_offset); + { + if (tgt->list[i].key != NULL) + devaddrs[i] = (void *) (tgt->list[i].key->tgt->tgt_start + + tgt->list[i].key->tgt_offset); + else + devaddrs[i] = NULL; + } acc_dev->openacc.exec_func (tgt_fn, mapnum, hostaddrs, devaddrs, async, dims, tgt); @@ -299,6 +304,8 @@ GOACC_data_start (int device, 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) || 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 00000000000..9f402973d3d --- /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 00000000000..bf17fc586b9 --- /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 00000000000..134e2d1cf29 --- /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 00000000000..276a1727b2e --- /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.17.1 --------------196CD9E12987FD434A40FEB8--