From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: (qmail 50562 invoked by alias); 19 Jun 2018 17:01:13 -0000 Mailing-List: contact gcc-patches-help@gcc.gnu.org; run by ezmlm Precedence: bulk List-Id: List-Archive: List-Post: List-Help: Sender: gcc-patches-owner@gcc.gnu.org Received: (qmail 50109 invoked by uid 89); 19 Jun 2018 17:00:56 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-25.2 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=tile 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; Tue, 19 Jun 2018 17:00:53 +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 1fVJzl-00015T-6u from Cesar_Philippidis@mentor.com ; Tue, 19 Jun 2018 10:00:41 -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; Tue, 19 Jun 2018 10:00:37 -0700 Subject: Re: [OpenACC] Update OpenACC data clause semantics to the 2.5 behavior - middle end From: Cesar Philippidis To: "gcc-patches@gcc.gnu.org" , Jakub Jelinek References: <7fa7637f-e7f5-d43d-13f1-706c77e8e957@codesourcery.com> <836f376f-513d-dd29-9133-6526bfb59866@codesourcery.com> Message-ID: <1f984431-affd-f63e-fe7e-f52cf1817f7f@codesourcery.com> Date: Tue, 19 Jun 2018 17:01:00 -0000 User-Agent: Mozilla/5.0 (X11; Linux x86_64; rv:52.0) Gecko/20100101 Thunderbird/52.8.0 MIME-Version: 1.0 In-Reply-To: <836f376f-513d-dd29-9133-6526bfb59866@codesourcery.com> Content-Type: multipart/mixed; boundary="------------29BFDAA29815941E13F0679D" X-ClientProxiedBy: svr-orw-mbx-01.mgc.mentorg.com (147.34.90.201) To svr-orw-mbx-01.mgc.mentorg.com (147.34.90.201) X-SW-Source: 2018-06/txt/msg01165.txt.bz2 --------------29BFDAA29815941E13F0679D Content-Type: text/plain; charset="utf-8" Content-Transfer-Encoding: 7bit Content-length: 107 This patch implements the OpenACC 2.5 data clause semantics in the middle end. Is it OK for trunk? Cesar --------------29BFDAA29815941E13F0679D Content-Type: text/x-patch; name="0006-gcc-middle-end.patch" Content-Transfer-Encoding: 7bit Content-Disposition: attachment; filename="0006-gcc-middle-end.patch" Content-length: 13697 2018-06-19 Chung-Lin Tang Thomas Schwinge Cesar Philippidis gcc/c-family/ * c-pragma.h (enum pragma_omp_clause): Add PRAGMA_OACC_CLAUSE_{FINALIZE,IF_PRESENT}. Remove PRAGMA_OACC_CLAUSE_PRESENT_OR_{COPY,COPYIN,COPYOUT,CREATE}. gcc/ * gimplify.c (gimplify_scan_omp_clauses): Add support for OMP_CLAUSE_{IF_PRESENT,FINALIZE}. (gimplify_adjust_omp_clauses): Likewise. (gimplify_oacc_declare_1): Add support for GOMP_MAP_RELEASE, remove support for GOMP_MAP_FORCE_{ALLOC,TO,FROM,TOFROM}. (gimplify_omp_target_update): Update handling of acc update and enter/exit data. * omp-low.c (install_var_field): Remove unused parameter base_pointers_restrict. (scan_sharing_clauses): Remove base_pointers_restrict parameter. Update call to install_var_field. Handle OMP_CLAUSE_{IF_PRESENT, FINALIZE} (omp_target_base_pointers_restrict_p): Delete. (scan_omp_target): Update call to scan_sharing_clauses. * tree-core.h (enum omp_clause_code): Add OMP_CLAUSE_{IF_PRESENT, FINALIZE}. * tree-nested.c (convert_nonlocal_omp_clauses): Handle OMP_CLAUSE_{IF_PRESENT,FINALIZE}. (convert_local_omp_clauses): Likewise. * tree-pretty-print.c (dump_omp_clause): Likewise. * tree.c (omp_clause_num_ops): Add entries for OMP_CLAUSE_{IF_PRESENT, FINALIZE}. (omp_clause_code_name): Likewise. >From f79b0e6f0d796dc18ef1faf20b9fad0b7feeaa94 Mon Sep 17 00:00:00 2001 From: Cesar Philippidis Date: Tue, 19 Jun 2018 09:32:20 -0700 Subject: [PATCH 6/7] gcc middle end --- gcc/c-family/c-pragma.h | 6 +-- gcc/gimplify.c | 67 ++++++++++++++++++++++------- gcc/omp-low.c | 93 +++++------------------------------------ gcc/tree-core.h | 8 +++- gcc/tree-nested.c | 4 ++ gcc/tree-pretty-print.c | 6 +++ gcc/tree.c | 8 +++- 7 files changed, 88 insertions(+), 104 deletions(-) diff --git a/gcc/c-family/c-pragma.h b/gcc/c-family/c-pragma.h index c70380c211b..b322547b11a 100644 --- a/gcc/c-family/c-pragma.h +++ b/gcc/c-family/c-pragma.h @@ -138,16 +138,13 @@ enum pragma_omp_clause { PRAGMA_OACC_CLAUSE_DELETE, PRAGMA_OACC_CLAUSE_DEVICEPTR, PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT, + PRAGMA_OACC_CLAUSE_FINALIZE, PRAGMA_OACC_CLAUSE_GANG, PRAGMA_OACC_CLAUSE_HOST, PRAGMA_OACC_CLAUSE_INDEPENDENT, PRAGMA_OACC_CLAUSE_NUM_GANGS, PRAGMA_OACC_CLAUSE_NUM_WORKERS, PRAGMA_OACC_CLAUSE_PRESENT, - PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY, - PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN, - PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT, - PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE, PRAGMA_OACC_CLAUSE_SELF, PRAGMA_OACC_CLAUSE_SEQ, PRAGMA_OACC_CLAUSE_TILE, @@ -156,6 +153,7 @@ enum pragma_omp_clause { PRAGMA_OACC_CLAUSE_VECTOR_LENGTH, PRAGMA_OACC_CLAUSE_WAIT, PRAGMA_OACC_CLAUSE_WORKER, + PRAGMA_OACC_CLAUSE_IF_PRESENT, PRAGMA_OACC_CLAUSE_COLLAPSE = PRAGMA_OMP_CLAUSE_COLLAPSE, PRAGMA_OACC_CLAUSE_COPYIN = PRAGMA_OMP_CLAUSE_COPYIN, PRAGMA_OACC_CLAUSE_DEVICE = PRAGMA_OMP_CLAUSE_DEVICE, diff --git a/gcc/gimplify.c b/gcc/gimplify.c index 1523a27e828..97543ed5f70 100644 --- a/gcc/gimplify.c +++ b/gcc/gimplify.c @@ -8524,6 +8524,8 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, case OMP_CLAUSE_NOGROUP: case OMP_CLAUSE_THREADS: case OMP_CLAUSE_SIMD: + case OMP_CLAUSE_IF_PRESENT: + case OMP_CLAUSE_FINALIZE: break; case OMP_CLAUSE_DEFAULTMAP: @@ -9305,6 +9307,8 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p, case OMP_CLAUSE_AUTO: case OMP_CLAUSE_SEQ: case OMP_CLAUSE_TILE: + case OMP_CLAUSE_IF_PRESENT: + case OMP_CLAUSE_FINALIZE: break; default: @@ -9361,21 +9365,7 @@ gimplify_oacc_declare_1 (tree clause) switch (kind) { case GOMP_MAP_ALLOC: - case GOMP_MAP_FORCE_ALLOC: - case GOMP_MAP_FORCE_TO: - new_op = GOMP_MAP_DELETE; - ret = true; - break; - - case GOMP_MAP_FORCE_FROM: - OMP_CLAUSE_SET_MAP_KIND (clause, GOMP_MAP_FORCE_ALLOC); - new_op = GOMP_MAP_FORCE_FROM; - ret = true; - break; - - case GOMP_MAP_FORCE_TOFROM: - OMP_CLAUSE_SET_MAP_KIND (clause, GOMP_MAP_FORCE_TO); - new_op = GOMP_MAP_FORCE_FROM; + new_op = GOMP_MAP_RELEASE; ret = true; break; @@ -10817,6 +10807,53 @@ gimplify_omp_target_update (tree *expr_p, gimple_seq *pre_p) ort, TREE_CODE (expr)); gimplify_adjust_omp_clauses (pre_p, NULL, &OMP_STANDALONE_CLAUSES (expr), TREE_CODE (expr)); + if (TREE_CODE (expr) == OACC_UPDATE + && omp_find_clause (OMP_STANDALONE_CLAUSES (expr), + OMP_CLAUSE_IF_PRESENT)) + { + /* The runtime uses GOMP_MAP_{TO,FROM} to denote the if_present + clause. */ + for (tree c = OMP_STANDALONE_CLAUSES (expr); c; c = OMP_CLAUSE_CHAIN (c)) + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP) + switch (OMP_CLAUSE_MAP_KIND (c)) + { + case GOMP_MAP_FORCE_TO: + OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_TO); + break; + case GOMP_MAP_FORCE_FROM: + OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_FROM); + break; + default: + break; + } + } + else if (TREE_CODE (expr) == OACC_EXIT_DATA + && omp_find_clause (OMP_STANDALONE_CLAUSES (expr), + OMP_CLAUSE_FINALIZE)) + { + /* Use GOMP_MAP_DELETE/GOMP_MAP_FORCE_FROM to denote that "finalize" + semantics apply to all mappings of this OpenACC directive. */ + bool finalize_marked = false; + for (tree c = OMP_STANDALONE_CLAUSES (expr); c; c = OMP_CLAUSE_CHAIN (c)) + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP) + switch (OMP_CLAUSE_MAP_KIND (c)) + { + case GOMP_MAP_FROM: + OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_FORCE_FROM); + finalize_marked = true; + break; + case GOMP_MAP_RELEASE: + OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_DELETE); + finalize_marked = true; + break; + default: + /* Check consistency: libgomp relies on the very first data + mapping clause being marked, so make sure we did that before + any other mapping clauses. */ + gcc_assert (finalize_marked); + break; + } + } stmt = gimple_build_omp_target (NULL, kind, OMP_STANDALONE_CLAUSES (expr)); gimplify_seq_add_stmt (pre_p, stmt); diff --git a/gcc/omp-low.c b/gcc/omp-low.c index ba6c705cf8b..c591231d8f1 100644 --- a/gcc/omp-low.c +++ b/gcc/omp-low.c @@ -642,8 +642,7 @@ build_sender_ref (tree var, omp_context *ctx) BASE_POINTERS_RESTRICT, declare the field with restrict. */ static void -install_var_field (tree var, bool by_ref, int mask, omp_context *ctx, - bool base_pointers_restrict = false) +install_var_field (tree var, bool by_ref, int mask, omp_context *ctx) { tree field, type, sfield = NULL_TREE; splay_tree_key key = (splay_tree_key) var; @@ -674,11 +673,7 @@ install_var_field (tree var, bool by_ref, int mask, omp_context *ctx, type = build_pointer_type (build_pointer_type (type)); } else if (by_ref) - { - type = build_pointer_type (type); - if (base_pointers_restrict) - type = build_qualified_type (type, TYPE_QUAL_RESTRICT); - } + type = build_pointer_type (type); else if ((mask & 3) == 1 && omp_is_reference (var)) type = TREE_TYPE (type); @@ -992,12 +987,10 @@ fixup_child_record_type (omp_context *ctx) } /* Instantiate decls as necessary in CTX to satisfy the data sharing - specified by CLAUSES. If BASE_POINTERS_RESTRICT, install var field with - restrict. */ + specified by CLAUSES. */ static void -scan_sharing_clauses (tree clauses, omp_context *ctx, - bool base_pointers_restrict = false) +scan_sharing_clauses (tree clauses, omp_context *ctx) { tree c, decl; bool scan_array_reductions = false; @@ -1256,8 +1249,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx, && TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE) install_var_field (decl, true, 7, ctx); else - install_var_field (decl, true, 3, ctx, - base_pointers_restrict); + install_var_field (decl, true, 3, ctx); if (is_gimple_omp_offloaded (ctx->stmt) && !OMP_CLAUSE_MAP_IN_REDUCTION (c)) install_var_local (decl, ctx); @@ -1328,6 +1320,8 @@ scan_sharing_clauses (tree clauses, omp_context *ctx, case OMP_CLAUSE_TILE: case OMP_CLAUSE__SIMT_: case OMP_CLAUSE_DEFAULT: + case OMP_CLAUSE_IF_PRESENT: + case OMP_CLAUSE_FINALIZE: break; case OMP_CLAUSE_ALIGNED: @@ -1499,6 +1493,8 @@ scan_sharing_clauses (tree clauses, omp_context *ctx, case OMP_CLAUSE_TILE: case OMP_CLAUSE__GRIDDIM_: case OMP_CLAUSE__SIMT_: + case OMP_CLAUSE_IF_PRESENT: + case OMP_CLAUSE_FINALIZE: break; case OMP_CLAUSE__CACHE_: @@ -2266,68 +2262,6 @@ scan_omp_single (gomp_single *stmt, omp_context *outer_ctx) layout_type (ctx->record_type); } -/* Return true if the CLAUSES of an omp target guarantee that the base pointers - used in the corresponding offloaded function are restrict. */ - -static bool -omp_target_base_pointers_restrict_p (tree clauses) -{ - /* The analysis relies on the GOMP_MAP_FORCE_* mapping kinds, which are only - used by OpenACC. */ - if (flag_openacc == 0) - return false; - - /* I. Basic example: - - void foo (void) - { - unsigned int a[2], b[2]; - - #pragma acc kernels \ - copyout (a) \ - copyout (b) - { - a[0] = 0; - b[0] = 1; - } - } - - After gimplification, we have: - - #pragma omp target oacc_kernels \ - map(force_from:a [len: 8]) \ - map(force_from:b [len: 8]) - { - a[0] = 0; - b[0] = 1; - } - - Because both mappings have the force prefix, we know that they will be - allocated when calling the corresponding offloaded function, which means we - can mark the base pointers for a and b in the offloaded function as - restrict. */ - - tree c; - for (c = clauses; c; c = OMP_CLAUSE_CHAIN (c)) - { - if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP) - return false; - - switch (OMP_CLAUSE_MAP_KIND (c)) - { - case GOMP_MAP_FORCE_ALLOC: - case GOMP_MAP_FORCE_TO: - case GOMP_MAP_FORCE_FROM: - case GOMP_MAP_FORCE_TOFROM: - break; - default: - return false; - } - } - - return true; -} - /* Scan a GIMPLE_OMP_TARGET. */ static void @@ -2349,20 +2283,13 @@ scan_omp_target (gomp_target *stmt, omp_context *outer_ctx) TYPE_NAME (ctx->record_type) = name; TYPE_ARTIFICIAL (ctx->record_type) = 1; - bool base_pointers_restrict = false; if (offloaded) { create_omp_child_function (ctx, false); gimple_omp_target_set_child_fn (stmt, ctx->cb.dst_fn); - - base_pointers_restrict = omp_target_base_pointers_restrict_p (clauses); - if (base_pointers_restrict - && dump_file && (dump_flags & TDF_DETAILS)) - fprintf (dump_file, - "Base pointers in offloaded function are restrict\n"); } - scan_sharing_clauses (clauses, ctx, base_pointers_restrict); + scan_sharing_clauses (clauses, ctx); scan_omp (gimple_omp_body_ptr (stmt), ctx); if (TYPE_FIELDS (ctx->record_type) == NULL) diff --git a/gcc/tree-core.h b/gcc/tree-core.h index 2bebb22a7e9..4a04e9e8b26 100644 --- a/gcc/tree-core.h +++ b/gcc/tree-core.h @@ -454,7 +454,13 @@ enum omp_clause_code { /* OpenMP internal-only clause to specify grid dimensions of a gridified kernel. */ - OMP_CLAUSE__GRIDDIM_ + OMP_CLAUSE__GRIDDIM_, + + /* OpenACC clause: if_present. */ + OMP_CLAUSE_IF_PRESENT, + + /* OpenACC clause: finalize. */ + OMP_CLAUSE_FINALIZE }; #undef DEFTREESTRUCT diff --git a/gcc/tree-nested.c b/gcc/tree-nested.c index b335d6b0afe..257ceae6f2d 100644 --- a/gcc/tree-nested.c +++ b/gcc/tree-nested.c @@ -1333,6 +1333,8 @@ convert_nonlocal_omp_clauses (tree *pclauses, struct walk_stmt_info *wi) case OMP_CLAUSE_SEQ: case OMP_CLAUSE_INDEPENDENT: case OMP_CLAUSE_AUTO: + case OMP_CLAUSE_IF_PRESENT: + case OMP_CLAUSE_FINALIZE: break; /* The following clause belongs to the OpenACC cache directive, which @@ -2022,6 +2024,8 @@ convert_local_omp_clauses (tree *pclauses, struct walk_stmt_info *wi) case OMP_CLAUSE_SEQ: case OMP_CLAUSE_INDEPENDENT: case OMP_CLAUSE_AUTO: + case OMP_CLAUSE_IF_PRESENT: + case OMP_CLAUSE_FINALIZE: break; /* The following clause belongs to the OpenACC cache directive, which diff --git a/gcc/tree-pretty-print.c b/gcc/tree-pretty-print.c index 63ec823c0ba..e65c40a41a3 100644 --- a/gcc/tree-pretty-print.c +++ b/gcc/tree-pretty-print.c @@ -1045,6 +1045,12 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags) false); pp_right_paren (pp); break; + case OMP_CLAUSE_IF_PRESENT: + pp_string (pp, "if_present"); + break; + case OMP_CLAUSE_FINALIZE: + pp_string (pp, "finalize"); + break; default: /* Should never happen. */ diff --git a/gcc/tree.c b/gcc/tree.c index 889d88c50b4..cf0b5f6dad3 100644 --- a/gcc/tree.c +++ b/gcc/tree.c @@ -343,6 +343,8 @@ unsigned const char omp_clause_num_ops[] = 1, /* OMP_CLAUSE_VECTOR_LENGTH */ 3, /* OMP_CLAUSE_TILE */ 2, /* OMP_CLAUSE__GRIDDIM_ */ + 0, /* OMP_CLAUSE_IF_PRESENT */ + 0, /* OMP_CLAUSE_FINALIZE */ }; const char * const omp_clause_code_name[] = @@ -413,7 +415,9 @@ const char * const omp_clause_code_name[] = "num_workers", "vector_length", "tile", - "_griddim_" + "_griddim_", + "if_present", + "finalize", }; @@ -11579,6 +11583,8 @@ walk_tree_1 (tree *tp, walk_tree_fn func, void *data, case OMP_CLAUSE_SEQ: case OMP_CLAUSE_TILE: case OMP_CLAUSE__SIMT_: + case OMP_CLAUSE_IF_PRESENT: + case OMP_CLAUSE_FINALIZE: WALK_SUBTREE_TAIL (OMP_CLAUSE_CHAIN (*tp)); case OMP_CLAUSE_LASTPRIVATE: -- 2.17.1 --------------29BFDAA29815941E13F0679D--