From patchwork Wed Dec 23 11:02:19 2015 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 8bit X-Patchwork-Submitter: Thomas Schwinge X-Patchwork-Id: 560449 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@bilbo.ozlabs.org Received: from sourceware.org (server1.sourceware.org [209.132.180.131]) (using TLSv1.2 with cipher ECDHE-RSA-AES256-GCM-SHA384 (256/256 bits)) (No client certificate requested) by ozlabs.org (Postfix) with ESMTPS id 31E37140BFD for ; Wed, 23 Dec 2015 22:02:48 +1100 (AEDT) Authentication-Results: ozlabs.org; dkim=pass (1024-bit key; unprotected) header.d=gcc.gnu.org header.i=@gcc.gnu.org header.b=Gxm8w8xe; dkim-atps=neutral DomainKey-Signature: a=rsa-sha1; c=nofws; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender:from :to:cc:subject:in-reply-to:references:date:message-id :mime-version:content-type; q=dns; s=default; b=aFCD7iIIB7viHPpx C2Tjo2RVN8dcYgybAUppzR9cVzRRriDa50C9U+bEUhM517jbSqyJwu6DjlaGSgDF 6HQSQRq7BHgtxOIzX4PC9DSqUQ+44ilxXy2PjqShIy6S0dgqCEdcw2JfSqRD/FZJ /F+Mm7/WRCIZmGqXw1NdzH1Hpd0= DKIM-Signature: v=1; a=rsa-sha1; c=relaxed; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender:from :to:cc:subject:in-reply-to:references:date:message-id :mime-version:content-type; s=default; bh=o8r6eXIbwNat6E9SMuX3ZD JLMa0=; b=Gxm8w8xeYj7UL+DqDL/1HyAwY03WdFuzEMuv+jrQwmMB5X/6R5CmZW N/oAea+XvUa/OWjKT8Vv2RMBD3/TGyNk640hR1mXCA2dRNuuWZW0GunZdNF/I8vQ 497YuHyfxn3m9IgMmIAc/CsN5MGga3WKsS/240QfjpPAZPKdPAUY8= Received: (qmail 99451 invoked by alias); 23 Dec 2015 11:02:40 -0000 Mailing-List: contact gcc-patches-help@gcc.gnu.org; run by ezmlm Precedence: bulk List-Id: List-Unsubscribe: List-Archive: List-Post: List-Help: Sender: gcc-patches-owner@gcc.gnu.org Delivered-To: mailing list gcc-patches@gcc.gnu.org Received: (qmail 99427 invoked by uid 89); 23 Dec 2015 11:02:38 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-2.1 required=5.0 tests=AWL, BAYES_00, RCVD_IN_DNSWL_LOW, SPF_PASS autolearn=ham version=3.3.2 spammy=Merge, ack, 116106, 11610, 6 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; Wed, 23 Dec 2015 11:02:36 +0000 Received: from svr-orw-fem-03.mgc.mentorg.com ([147.34.97.39]) by relay1.mentorg.com with esmtp id 1aBhBh-0005YD-7B from Thomas_Schwinge@mentor.com ; Wed, 23 Dec 2015 03:02:33 -0800 Received: from tftp-cs (147.34.91.1) by svr-orw-fem-03.mgc.mentorg.com (147.34.97.39) with Microsoft SMTP Server id 14.3.224.2; Wed, 23 Dec 2015 03:02:32 -0800 Received: by tftp-cs (Postfix, from userid 49978) id 2567DC23E3; Wed, 23 Dec 2015 03:02:32 -0800 (PST) From: Thomas Schwinge To: , Julian Brown CC: "Joseph S. Myers" , Nathan Sidwell , Jakub Jelinek , James Norris Subject: Re: [Bulk] [OpenACC 0/7] host_data construct In-Reply-To: <20151026183422.GW478@tucnak.redhat.com> References: <56293476.5020801@codesourcery.com> <562A578E.4080907@codesourcery.com> <20151026183422.GW478@tucnak.redhat.com> User-Agent: Notmuch/0.9-125-g4686d11 (http://notmuchmail.org) Emacs/24.5.1 (i586-pc-linux-gnu) Date: Wed, 23 Dec 2015 12:02:19 +0100 Message-ID: <87twn9h1hg.fsf@kepler.schwinge.homeip.net> MIME-Version: 1.0 Hi! On Mon, 26 Oct 2015 19:34:22 +0100, Jakub Jelinek wrote: > Your use_device sounds very similar to use_device_ptr clause in OpenMP, > which is allowed on #pragma omp target data construct and is implemented > quite a bit differently from this; it is unclear if the OpenACC standard > requires this kind of implementation, or you just chose to implement it this > way. In particular, the GOMP_target_data call puts the variables mentioned > in the use_device_ptr clauses into the mapping structures (similarly how > map clause appears) and the corresponding vars are privatized within the > target data region (which is a host region, basically a fancy { } braces), > where the private variables contain the offloading device's pointers. ACK. As the OpenACC use_device clause implementation now completely matches the OpenMP use_device_ptr clause implementation, there is no use anymore for a separate OMP_CLAUSE_USE_DEVICE, so in r231926 I cleaned that up, as obvious: commit 9d5fd7c608fef6e7a9efbfc940545d49452c4e01 Author: tschwinge Date: Wed Dec 23 11:01:18 2015 +0000 Merge OMP_CLAUSE_USE_DEVICE into OMP_CLAUSE_USE_DEVICE_PTR gcc/c/ * c-parser.c (c_parser_oacc_clause_use_device): Merge function into... (c_parser_omp_clause_use_device_ptr): ... this function. Adjust all users. gcc/ * tree-core.h (enum omp_clause_code): Merge OMP_CLAUSE_USE_DEVICE into OMP_CLAUSE_USE_DEVICE_PTR. Adjust all users. git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/trunk@231926 138bc75d-0d04-0410-961f-82ee72b054a4 --- gcc/ChangeLog | 5 +++++ gcc/c/ChangeLog | 7 +++++++ gcc/c/c-parser.c | 16 +++++----------- gcc/c/c-typeck.c | 1 - gcc/cp/parser.c | 2 +- gcc/cp/pt.c | 2 -- gcc/cp/semantics.c | 1 - gcc/fortran/trans-openmp.c | 2 +- gcc/gimplify.c | 2 -- gcc/omp-low.c | 11 ++--------- gcc/testsuite/gfortran.dg/goacc/host_data-tree.f95 | 2 +- gcc/tree-core.h | 6 ++---- gcc/tree-nested.c | 2 -- gcc/tree-pretty-print.c | 3 --- gcc/tree.c | 3 --- 15 files changed, 24 insertions(+), 41 deletions(-) Grüße Thomas diff --git gcc/ChangeLog gcc/ChangeLog index aa28d10..d67b9c6 100644 --- gcc/ChangeLog +++ gcc/ChangeLog @@ -1,3 +1,8 @@ +2015-12-23 Thomas Schwinge + + * tree-core.h (enum omp_clause_code): Merge OMP_CLAUSE_USE_DEVICE + into OMP_CLAUSE_USE_DEVICE_PTR. Adjust all users. + 2015-12-23 David Sherwood * config/arm/iterators.md (VMAXMINFNM): New int iterator. diff --git gcc/c/ChangeLog gcc/c/ChangeLog index f99f426..7b275d8 100644 --- gcc/c/ChangeLog +++ gcc/c/ChangeLog @@ -1,3 +1,10 @@ +2015-12-23 Thomas Schwinge + + * c-parser.c (c_parser_oacc_clause_use_device): Merge function + into... + (c_parser_omp_clause_use_device_ptr): ... this function. Adjust + all users. + 2015-12-22 Marek Polacek PR c/69002 diff --git gcc/c/c-parser.c gcc/c/c-parser.c index 353e3da..8e754d0 100644 --- gcc/c/c-parser.c +++ gcc/c/c-parser.c @@ -11395,7 +11395,10 @@ c_parser_omp_clause_defaultmap (c_parser *parser, tree list) return list; } -/* OpenMP 4.5: +/* OpenACC 2.0: + use_device ( variable-list ) + + OpenMP 4.5: use_device_ptr ( variable-list ) */ static tree @@ -11730,15 +11733,6 @@ c_parser_oacc_clause_tile (c_parser *parser, tree list) return c; } -/* OpenACC 2.0: - use_device ( variable-list ) */ - -static tree -c_parser_oacc_clause_use_device (c_parser *parser, tree list) -{ - return c_parser_omp_var_list_parens (parser, OMP_CLAUSE_USE_DEVICE, list); -} - /* OpenACC: wait ( int-expr-list ) */ @@ -13058,7 +13052,7 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask, c_name = "tile"; break; case PRAGMA_OACC_CLAUSE_USE_DEVICE: - clauses = c_parser_oacc_clause_use_device (parser, clauses); + clauses = c_parser_omp_clause_use_device_ptr (parser, clauses); c_name = "use_device"; break; case PRAGMA_OACC_CLAUSE_VECTOR: diff --git gcc/c/c-typeck.c gcc/c/c-typeck.c index 928fcd5..7406bd4 100644 --- gcc/c/c-typeck.c +++ gcc/c/c-typeck.c @@ -13125,7 +13125,6 @@ c_finish_omp_clauses (tree clauses, bool is_omp, bool declare_simd) bitmap_set_bit (&map_head, DECL_UID (t)); goto check_dup_generic; - case OMP_CLAUSE_USE_DEVICE: case OMP_CLAUSE_IS_DEVICE_PTR: case OMP_CLAUSE_USE_DEVICE_PTR: t = OMP_CLAUSE_DECL (c); diff --git gcc/cp/parser.c gcc/cp/parser.c index 842dded..4829a77 100644 --- gcc/cp/parser.c +++ gcc/cp/parser.c @@ -32097,7 +32097,7 @@ cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask, c_name = "tile"; break; case PRAGMA_OACC_CLAUSE_USE_DEVICE: - clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE_USE_DEVICE, + clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE_USE_DEVICE_PTR, clauses); c_name = "use_device"; break; diff --git gcc/cp/pt.c gcc/cp/pt.c index dab15bd..4555b32 100644 --- gcc/cp/pt.c +++ gcc/cp/pt.c @@ -14425,7 +14425,6 @@ tsubst_omp_clauses (tree clauses, bool declare_simd, bool allow_fields, case OMP_CLAUSE_FROM: case OMP_CLAUSE_TO: case OMP_CLAUSE_MAP: - case OMP_CLAUSE_USE_DEVICE: case OMP_CLAUSE_USE_DEVICE_PTR: case OMP_CLAUSE_IS_DEVICE_PTR: OMP_CLAUSE_DECL (nc) @@ -14552,7 +14551,6 @@ tsubst_omp_clauses (tree clauses, bool declare_simd, bool allow_fields, case OMP_CLAUSE_COPYPRIVATE: case OMP_CLAUSE_LINEAR: case OMP_CLAUSE_REDUCTION: - case OMP_CLAUSE_USE_DEVICE: case OMP_CLAUSE_USE_DEVICE_PTR: case OMP_CLAUSE_IS_DEVICE_PTR: /* tsubst_expr on SCOPE_REF results in returning diff --git gcc/cp/semantics.c gcc/cp/semantics.c index ab9989a..37bf050 100644 --- gcc/cp/semantics.c +++ gcc/cp/semantics.c @@ -6886,7 +6886,6 @@ finish_omp_clauses (tree clauses, bool allow_fields, bool declare_simd) } break; - case OMP_CLAUSE_USE_DEVICE: case OMP_CLAUSE_IS_DEVICE_PTR: case OMP_CLAUSE_USE_DEVICE_PTR: field_ok = allow_fields; diff --git gcc/fortran/trans-openmp.c gcc/fortran/trans-openmp.c index 227964c..70a7722 100644 --- gcc/fortran/trans-openmp.c +++ gcc/fortran/trans-openmp.c @@ -1771,7 +1771,7 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, clause_code = OMP_CLAUSE_UNIFORM; goto add_clause; case OMP_LIST_USE_DEVICE: - clause_code = OMP_CLAUSE_USE_DEVICE; + clause_code = OMP_CLAUSE_USE_DEVICE_PTR; goto add_clause; case OMP_LIST_DEVICE_RESIDENT: clause_code = OMP_CLAUSE_DEVICE_RESIDENT; diff --git gcc/gimplify.c gcc/gimplify.c index 62b0e64..bc90401 100644 --- gcc/gimplify.c +++ gcc/gimplify.c @@ -7139,7 +7139,6 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, } goto do_notice; - case OMP_CLAUSE_USE_DEVICE: case OMP_CLAUSE_USE_DEVICE_PTR: flags = GOVD_FIRSTPRIVATE | GOVD_EXPLICIT; goto do_add; @@ -8051,7 +8050,6 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p, case OMP_CLAUSE_ASYNC: case OMP_CLAUSE_WAIT: case OMP_CLAUSE_DEVICE_RESIDENT: - case OMP_CLAUSE_USE_DEVICE: case OMP_CLAUSE_INDEPENDENT: case OMP_CLAUSE_NUM_GANGS: case OMP_CLAUSE_NUM_WORKERS: diff --git gcc/omp-low.c gcc/omp-low.c index 676b1df..a0c3e1c 100644 --- gcc/omp-low.c +++ gcc/omp-low.c @@ -1957,7 +1957,6 @@ scan_sharing_clauses (tree clauses, omp_context *ctx, install_var_local (decl, ctx); break; - case OMP_CLAUSE_USE_DEVICE: case OMP_CLAUSE_USE_DEVICE_PTR: decl = OMP_CLAUSE_DECL (c); if (TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE) @@ -2314,7 +2313,6 @@ scan_sharing_clauses (tree clauses, omp_context *ctx, case OMP_CLAUSE_SIMD: case OMP_CLAUSE_NOGROUP: case OMP_CLAUSE_DEFAULTMAP: - case OMP_CLAUSE_USE_DEVICE: case OMP_CLAUSE_USE_DEVICE_PTR: case OMP_CLAUSE__CILK_FOR_COUNT_: case OMP_CLAUSE_ASYNC: @@ -15288,7 +15286,6 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) } break; - case OMP_CLAUSE_USE_DEVICE: case OMP_CLAUSE_USE_DEVICE_PTR: case OMP_CLAUSE_IS_DEVICE_PTR: var = OMP_CLAUSE_DECL (c); @@ -15674,14 +15671,12 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) build_int_cstu (tkind_type, tkind)); break; - case OMP_CLAUSE_USE_DEVICE: case OMP_CLAUSE_USE_DEVICE_PTR: case OMP_CLAUSE_IS_DEVICE_PTR: ovar = OMP_CLAUSE_DECL (c); var = lookup_decl_in_outer_ctx (ovar, ctx); x = build_sender_ref (ovar, ctx); - if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_PTR - || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE) + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_PTR) tkind = GOMP_MAP_USE_DEVICE_PTR; else tkind = GOMP_MAP_FIRSTPRIVATE_INT; @@ -15884,12 +15879,10 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) gimple_build_assign (new_var, x)); } break; - case OMP_CLAUSE_USE_DEVICE: case OMP_CLAUSE_USE_DEVICE_PTR: case OMP_CLAUSE_IS_DEVICE_PTR: var = OMP_CLAUSE_DECL (c); - if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_PTR - || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE) + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_PTR) x = build_sender_ref (var, ctx); else x = build_receiver_ref (var, false, ctx); diff --git gcc/testsuite/gfortran.dg/goacc/host_data-tree.f95 gcc/testsuite/gfortran.dg/goacc/host_data-tree.f95 index 7a5eea6..23aba8c 100644 --- gcc/testsuite/gfortran.dg/goacc/host_data-tree.f95 +++ gcc/testsuite/gfortran.dg/goacc/host_data-tree.f95 @@ -8,4 +8,4 @@ program test !$acc host_data use_device(i) !$acc end host_data end program test -! { dg-final { scan-tree-dump-times "pragma acc host_data use_device\\(i\\)" 1 "original" } } +! { dg-final { scan-tree-dump-times "pragma acc host_data use_device_ptr\\(i\\)" 1 "original" } } diff --git gcc/tree-core.h gcc/tree-core.h index 9cc64d9..5371378 100644 --- gcc/tree-core.h +++ gcc/tree-core.h @@ -302,7 +302,8 @@ enum omp_clause_code { OpenMP clause: map ({alloc:,to:,from:,tofrom:,}variable-list). */ OMP_CLAUSE_MAP, - /* OpenMP clause: use_device_ptr (variable-list). */ + /* OpenACC clause: use_device (variable_list). + OpenMP clause: use_device_ptr (variable-list). */ OMP_CLAUSE_USE_DEVICE_PTR, /* OpenMP clause: is_device_ptr (variable-list). */ @@ -315,9 +316,6 @@ enum omp_clause_code { /* OpenACC clause: device_resident (variable_list). */ OMP_CLAUSE_DEVICE_RESIDENT, - /* OpenACC clause: use_device (variable_list). */ - OMP_CLAUSE_USE_DEVICE, - /* OpenACC clause: gang [(gang-argument-list)]. Where gang-argument-list: [gang-argument-list, ] gang-argument diff --git gcc/tree-nested.c gcc/tree-nested.c index 3a9479a..8211a12 100644 --- gcc/tree-nested.c +++ gcc/tree-nested.c @@ -1072,7 +1072,6 @@ convert_nonlocal_omp_clauses (tree *pclauses, struct walk_stmt_info *wi) case OMP_CLAUSE_SHARED: case OMP_CLAUSE_TO_DECLARE: case OMP_CLAUSE_LINK: - case OMP_CLAUSE_USE_DEVICE: case OMP_CLAUSE_USE_DEVICE_PTR: case OMP_CLAUSE_IS_DEVICE_PTR: do_decl_clause: @@ -1744,7 +1743,6 @@ convert_local_omp_clauses (tree *pclauses, struct walk_stmt_info *wi) case OMP_CLAUSE_SHARED: case OMP_CLAUSE_TO_DECLARE: case OMP_CLAUSE_LINK: - case OMP_CLAUSE_USE_DEVICE: case OMP_CLAUSE_USE_DEVICE_PTR: case OMP_CLAUSE_IS_DEVICE_PTR: do_decl_clause: diff --git gcc/tree-pretty-print.c gcc/tree-pretty-print.c index caec760..c4a180e 100644 --- gcc/tree-pretty-print.c +++ gcc/tree-pretty-print.c @@ -327,9 +327,6 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, int flags) case OMP_CLAUSE_DEVICE_RESIDENT: name = "device_resident"; goto print_remap; - case OMP_CLAUSE_USE_DEVICE: - name = "use_device"; - goto print_remap; case OMP_CLAUSE_TO_DECLARE: name = "to"; goto print_remap; diff --git gcc/tree.c gcc/tree.c index 2190cae..d837609 100644 --- gcc/tree.c +++ gcc/tree.c @@ -282,7 +282,6 @@ unsigned const char omp_clause_num_ops[] = 1, /* OMP_CLAUSE_IS_DEVICE_PTR */ 2, /* OMP_CLAUSE__CACHE_ */ 1, /* OMP_CLAUSE_DEVICE_RESIDENT */ - 1, /* OMP_CLAUSE_USE_DEVICE */ 2, /* OMP_CLAUSE_GANG */ 1, /* OMP_CLAUSE_ASYNC */ 1, /* OMP_CLAUSE_WAIT */ @@ -354,7 +353,6 @@ const char * const omp_clause_code_name[] = "is_device_ptr", "_cache_", "device_resident", - "use_device", "gang", "async", "wait", @@ -11612,7 +11610,6 @@ walk_tree_1 (tree *tp, walk_tree_fn func, void *data, /* FALLTHRU */ case OMP_CLAUSE_DEVICE_RESIDENT: - case OMP_CLAUSE_USE_DEVICE: case OMP_CLAUSE_ASYNC: case OMP_CLAUSE_WAIT: case OMP_CLAUSE_WORKER: