From patchwork Thu Nov 12 01:07:58 2015 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 8bit X-Patchwork-Submitter: James Norris X-Patchwork-Id: 543162 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 397A2140549 for ; Thu, 12 Nov 2015 12:08:22 +1100 (AEDT) Authentication-Results: ozlabs.org; dkim=pass (1024-bit key; unprotected) header.d=gcc.gnu.org header.i=@gcc.gnu.org header.b=f2ihjEjj; 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 :message-id:date:from:mime-version:to:cc:subject:references :in-reply-to:content-type; q=dns; s=default; b=UQicZaZ2VdIEEiD/u GvA3gZoSK7LBlKc0Kt81SWNDfeOPNb473fWHT30RswLtVgtCaFHkUc+dE0hDWJ0O mR0a8Ayh1CSH2FxHRMEiDXsmb2r6M2iNmfGIkyC86c3nIT+vThzOFEJUY5iGIX5a 2lrqAcGytc17gkACt393eGGAPg= 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 :message-id:date:from:mime-version:to:cc:subject:references :in-reply-to:content-type; s=default; bh=Ci9Vx3DNcYpVFTA0w5pYJ3v CbBk=; b=f2ihjEjjwpcGMagNXcJQOG2Ef4Qn0zFpakjYRJSO5Ef3vKEBVL3gI02 8Ov7jd0c/LzAPTfxiU5pdPvSAvZarpZAYT4zHi1BCRsL/RPNZF3S/v353XO7xvZ0 LJBxWfhFkH0kp27KZpyVDA8fwwKo7FUqrwtVxH/w1s8fBXvCq2Pc= Received: (qmail 66449 invoked by alias); 12 Nov 2015 01:08:09 -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 66429 invoked by uid 89); 12 Nov 2015 01:08:08 -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 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; Thu, 12 Nov 2015 01:08:04 +0000 Received: from svr-orw-fem-02x.mgc.mentorg.com ([147.34.96.206] helo=SVR-ORW-FEM-02.mgc.mentorg.com) by relay1.mentorg.com with esmtp id 1ZwgMq-0006is-JX from James_Norris@mentor.com ; Wed, 11 Nov 2015 17:08:00 -0800 Received: from [172.30.80.60] (147.34.91.1) by svr-orw-fem-02.mgc.mentorg.com (147.34.96.168) with Microsoft SMTP Server id 14.3.224.2; Wed, 11 Nov 2015 17:07:59 -0800 Message-ID: <5643E66E.9000202@codesourcery.com> Date: Wed, 11 Nov 2015 19:07:58 -0600 From: James Norris User-Agent: Mozilla/5.0 (X11; Linux x86_64; rv:31.0) Gecko/20100101 Thunderbird/31.7.0 MIME-Version: 1.0 To: Jakub Jelinek , James Norris CC: Thomas Schwinge , "Joseph S. Myers" , GCC Patches Subject: Re: [OpenACC] declare directive References: <562FDBF8.1040105@mentor.com> <5638E164.5010207@codesourcery.com> <87611h1zi7.fsf@kepler.schwinge.homeip.net> <563CD07A.3000404@codesourcery.com> <20151106190352.GG5675@tucnak.redhat.com> <563F6BA5.9020606@codesourcery.com> <5640C35C.9030907@codesourcery.com> <20151109162117.GU5675@tucnak.redhat.com> <56412830.9030100@codesourcery.com> <20151111083233.GE5675@tucnak.redhat.com> In-Reply-To: <20151111083233.GE5675@tucnak.redhat.com> X-TagToolbar-Keys: D20151111190758269 Jakub, The attached patch and ChangeLog reflect the updates from your review: https://gcc.gnu.org/ml/gcc-patches/2015-11/msg01317.html. Highlights.... The following issue was handled by Dominique d'Humières in: https://gcc.gnu.org/ml/gcc-patches/2015-11/msg01375.html On 11/11/2015 02:32 AM, Jakub Jelinek wrote: > On Mon, Nov 09, 2015 at 05:11:44PM -0600, James Norris wrote: >> >diff --git a/gcc/c-family/c-pragma.h b/gcc/c-family/c-pragma.h >> >index 953c4e3..c6a2981 100644 >> >--- a/gcc/c-family/c-pragma.h >> >+++ b/gcc/c-family/c-pragma.h >> >@@ -30,6 +30,7 @@ enum pragma_kind { >> > PRAGMA_OACC_ATOMIC, >> > PRAGMA_OACC_CACHE, >> > PRAGMA_OACC_DATA, >> >+ PRAGMA_OACC_DECLARE, >> > PRAGMA_OACC_ENTER_DATA, >> > PRAGMA_OACC_EXIT_DATA, >> > PRAGMA_OACC_KERNELS, > This change will make PR68271 even worse, so would be really nice to > get that fixed first. With the addition of: https://gcc.gnu.org/ml/gcc-patches/2015-11/msg01372.html, additional conditions were added to the following as you called out in your review of: https://gcc.gnu.org/ml/gcc-patches/2015-11/msg00703.html. On 11/06/2015 01:03 PM, Jakub Jelinek wrote: >> @@ -5841,6 +5863,8 @@ omp_default_clause (struct gimplify_omp_ctx *ctx, tree decl, >> flags |= GOVD_FIRSTPRIVATE; >> break; >> case OMP_CLAUSE_DEFAULT_UNSPECIFIED: >> + if (is_global_var (decl) && device_resident_p (decl)) >> + flags |= GOVD_MAP_TO_ONLY | GOVD_MAP; > > I don't think you want to do this except for (selected or all?) > OpenACC contexts. Say, I don't see why one couldn't e.g. try to mix > OpenMP host parallelization or tasking with OpenACC offloading, > and that affecting in weird way OpenMP semantics. > With the addition of routine directive support, additional run-time tests were added. OK? Thanks, Jim 2015-XX-XX James Norris Joseph Myers gcc/c-family/ * c-pragma.c (oacc_pragmas): Add entry for declare directive. * c-pragma.h (enum pragma_kind): Add PRAGMA_OACC_DECLARE. (enum pragma_omp_clause): Add PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT and PRAGMA_OACC_CLAUSE_LINK. gcc/c/ * c-parser.c (c_parser_pragma): Handle PRAGMA_OACC_DECLARE. (c_parser_omp_clause_name): Handle 'device_resident' clause. (c_parser_oacc_data_clause): Handle PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT and PRAGMA_OMP_CLAUSE_LINK. (c_parser_oacc_all_clauses): Handle PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT and PRAGMA_OACC_CLAUSE_LINK. (OACC_DECLARE_CLAUSE_MASK): New definition. (c_parser_oacc_declare): New function. gcc/cp/ * parser.c (cp_parser_omp_clause_name): Handle 'device_resident' clause. (cp_parser_oacc_data_clause): Handle PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT and PRAGMA_OMP_CLAUSE_LINK. (cp_paser_oacc_all_clauses): Handle PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT and PRAGMA_OMP_CLAUSE_LINK. (OACC_DECLARE_CLAUSE_MASK): New definition. (cp_parser_oacc_declare): New function. (cp_parser_pragma): Handle PRAGMA_OACC_DECLARE. * pt.c (tsubst_expr): Handle OACC_DECLARE. gcc/ * gimple-pretty-print.c (dump_gimple_omp_target): Handle GF_OMP_TARGET_KIND_OACC_DECLARE. * gimple.h (enum gf_mask): Add GF_OMP_TARGET_KIND_OACC_DECLARE. (is_gomple_omp_oacc): Handle GF_OMP_TARGET_KIND_OACC_DECLARE. * gimplify.c (oacc_declare_returns): New. (gimplify_bind_expr): Prepend 'exit' stmt to cleanup. (device_resident_p): New function. (omp_default_clause): Handle device_resident clause. (gimplify_oacc_declare_1, gimplify_oacc_declare): New functions. (gimplify_expr): Handle OACC_DECLARE. * omp-builtins.def (BUILT_IN_GOACC_DECLARE): New builtin. * omp-low.c (expand_omp_target): Handle GF_OMP_TARGET_KIND_OACC_DECLARE and BUILTIN_GOACC_DECLARE. (build_omp_regions_1): Handlde GF_OMP_TARGET_KIND_OACC_DECLARE. (lower_omp_target): Handle GF_OMP_TARGET_KIND_OACC_DECLARE, GOMP_MAP_DEVICE_RESIDENT and GOMP_MAP_LINK. (make_gimple_omp_edges): Handle GF_OMP_TARGET_KIND_OACC_DECLARE. * tree-pretty-print.c (dump_omp_clause): Handle GOMP_MAP_LINK and GOMP_MAP_DEVICE_RESIDENT. gcc/testsuite * c-c++-common/goacc/declare-1.c: New test. * c-c++-common/goacc/declare-2.c: Likewise. include/ * gomp-constants.h (enum gomp_map_kind): Add GOMP_MAP_DEVICE_RESIDENT and GOMP_MAP_LINK. libgomp/ * libgomp.map (GOACC_2.0.1): Export GOACC_declare. * oacc-parallel.c (GOACC_declare): New function. * testsuite/libgomp.oacc-c-c++-common/declare-1.c: New test. * testsuite/libgomp.oacc-c-c++-common/declare-2.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/declare-4.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/declare-5.c: Likewise. * testsuite/libgomp.oacc-c++/declare-1.C: Likewise. diff --git a/gcc/c-family/c-pragma.c b/gcc/c-family/c-pragma.c index f86ed38..12c3e75 100644 --- a/gcc/c-family/c-pragma.c +++ b/gcc/c-family/c-pragma.c @@ -1248,6 +1248,7 @@ static const struct omp_pragma_def oacc_pragmas[] = { { "atomic", PRAGMA_OACC_ATOMIC }, { "cache", PRAGMA_OACC_CACHE }, { "data", PRAGMA_OACC_DATA }, + { "declare", PRAGMA_OACC_DECLARE }, { "enter", PRAGMA_OACC_ENTER_DATA }, { "exit", PRAGMA_OACC_EXIT_DATA }, { "kernels", PRAGMA_OACC_KERNELS }, diff --git a/gcc/c-family/c-pragma.h b/gcc/c-family/c-pragma.h index afeceff..999ac67 100644 --- a/gcc/c-family/c-pragma.h +++ b/gcc/c-family/c-pragma.h @@ -30,6 +30,7 @@ enum pragma_kind { PRAGMA_OACC_ATOMIC, PRAGMA_OACC_CACHE, PRAGMA_OACC_DATA, + PRAGMA_OACC_DECLARE, PRAGMA_OACC_ENTER_DATA, PRAGMA_OACC_EXIT_DATA, PRAGMA_OACC_KERNELS, @@ -152,6 +153,7 @@ enum pragma_omp_clause { PRAGMA_OACC_CLAUSE_CREATE, PRAGMA_OACC_CLAUSE_DELETE, PRAGMA_OACC_CLAUSE_DEVICEPTR, + PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT, PRAGMA_OACC_CLAUSE_GANG, PRAGMA_OACC_CLAUSE_HOST, PRAGMA_OACC_CLAUSE_INDEPENDENT, @@ -176,7 +178,8 @@ enum pragma_omp_clause { PRAGMA_OACC_CLAUSE_FIRSTPRIVATE = PRAGMA_OMP_CLAUSE_FIRSTPRIVATE, PRAGMA_OACC_CLAUSE_IF = PRAGMA_OMP_CLAUSE_IF, PRAGMA_OACC_CLAUSE_PRIVATE = PRAGMA_OMP_CLAUSE_PRIVATE, - PRAGMA_OACC_CLAUSE_REDUCTION = PRAGMA_OMP_CLAUSE_REDUCTION + PRAGMA_OACC_CLAUSE_REDUCTION = PRAGMA_OMP_CLAUSE_REDUCTION, + PRAGMA_OACC_CLAUSE_LINK = PRAGMA_OMP_CLAUSE_LINK }; extern struct cpp_reader* parse_in; diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c index 2484b92..b7806bd 100644 --- a/gcc/c/c-parser.c +++ b/gcc/c/c-parser.c @@ -1228,6 +1228,7 @@ static vec *c_parser_expr_list (c_parser *, bool, bool, vec **, location_t *, tree *, vec *, unsigned int * = NULL); +static void c_parser_oacc_declare (c_parser *); static void c_parser_oacc_enter_exit_data (c_parser *, bool); static void c_parser_oacc_update (c_parser *); static void c_parser_omp_construct (c_parser *); @@ -9714,6 +9715,10 @@ c_parser_pragma (c_parser *parser, enum pragma_context context) switch (id) { + case PRAGMA_OACC_DECLARE: + c_parser_oacc_declare (parser); + return false; + case PRAGMA_OACC_ENTER_DATA: c_parser_oacc_enter_exit_data (parser, true); return false; @@ -10003,6 +10008,8 @@ c_parser_omp_clause_name (c_parser *parser) result = PRAGMA_OMP_CLAUSE_DEVICE; else if (!strcmp ("deviceptr", p)) result = PRAGMA_OACC_CLAUSE_DEVICEPTR; + else if (!strcmp ("device_resident", p)) + result = PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT; else if (!strcmp ("dist_schedule", p)) result = PRAGMA_OMP_CLAUSE_DIST_SCHEDULE; break; @@ -10439,10 +10446,16 @@ c_parser_oacc_data_clause (c_parser *parser, pragma_omp_clause c_kind, case PRAGMA_OACC_CLAUSE_DEVICE: kind = GOMP_MAP_FORCE_TO; break; + case PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT: + kind = GOMP_MAP_DEVICE_RESIDENT; + break; case PRAGMA_OACC_CLAUSE_HOST: case PRAGMA_OACC_CLAUSE_SELF: kind = GOMP_MAP_FORCE_FROM; break; + case PRAGMA_OACC_CLAUSE_LINK: + kind = GOMP_MAP_LINK; + break; case PRAGMA_OACC_CLAUSE_PRESENT: kind = GOMP_MAP_FORCE_PRESENT; break; @@ -12724,6 +12737,10 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask, clauses = c_parser_oacc_data_clause_deviceptr (parser, clauses); c_name = "deviceptr"; break; + case PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT: + clauses = c_parser_oacc_data_clause (parser, c_kind, clauses); + c_name = "device_resident"; + break; case PRAGMA_OACC_CLAUSE_FIRSTPRIVATE: clauses = c_parser_omp_clause_firstprivate (parser, clauses); c_name = "firstprivate"; @@ -12746,6 +12763,10 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask, clauses); c_name = "independent"; break; + case PRAGMA_OACC_CLAUSE_LINK: + clauses = c_parser_oacc_data_clause (parser, c_kind, clauses); + c_name = "link"; + break; case PRAGMA_OACC_CLAUSE_NUM_GANGS: clauses = c_parser_omp_clause_num_gangs (parser, clauses); c_name = "num_gangs"; @@ -13203,6 +13224,161 @@ c_parser_oacc_data (location_t loc, c_parser *parser) } /* OpenACC 2.0: + # pragma acc declare oacc-data-clause[optseq] new-line +*/ + +#define OACC_DECLARE_CLAUSE_MASK \ + ( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPY) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYIN) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYOUT) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_CREATE) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICEPTR) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_LINK) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE) ) + +static void +c_parser_oacc_declare (c_parser *parser) +{ + location_t pragma_loc = c_parser_peek_token (parser)->location; + tree clauses, stmt, t, decl; + + bool error = false; + + c_parser_consume_pragma (parser); + + clauses = c_parser_oacc_all_clauses (parser, OACC_DECLARE_CLAUSE_MASK, + "#pragma acc declare"); + if (!clauses) + { + error_at (pragma_loc, + "no valid clauses specified in %<#pragma acc declare%>"); + return; + } + + for (t = clauses; t; t = OMP_CLAUSE_CHAIN (t)) + { + location_t loc = OMP_CLAUSE_LOCATION (t); + decl = OMP_CLAUSE_DECL (t); + if (!DECL_P (decl)) + { + error_at (loc, "array section in %<#pragma acc declare%>"); + error = true; + continue; + } + + switch (OMP_CLAUSE_MAP_KIND (t)) + { + case GOMP_MAP_FORCE_ALLOC: + case GOMP_MAP_FORCE_TO: + case GOMP_MAP_FORCE_DEVICEPTR: + case GOMP_MAP_DEVICE_RESIDENT: + break; + + case GOMP_MAP_POINTER: + /* Generated by c_finish_omp_clauses from array sections; + avoid spurious diagnostics. */ + break; + + case GOMP_MAP_LINK: + if (!global_bindings_p () + && (TREE_STATIC (decl) + || !DECL_EXTERNAL (decl))) + { + error_at (loc, + "%qD must be a global variable in" + "%<#pragma acc declare link%>", + decl); + error = true; + continue; + } + break; + + default: + if (global_bindings_p ()) + { + error_at (loc, "invalid OpenACC clause at file scope"); + error = true; + continue; + } + if (DECL_EXTERNAL (decl)) + { + error_at (loc, + "invalid use of % variable %qD " + "in %<#pragma acc declare%>", decl); + error = true; + continue; + } + else if (TREE_PUBLIC (decl)) + { + error_at (loc, + "invalid use of % variable %qD " + "in %<#pragma acc declare%>", decl); + error = true; + continue; + } + break; + } + + if (lookup_attribute ("omp declare target", DECL_ATTRIBUTES (decl)) + || lookup_attribute ("omp declare target link", + DECL_ATTRIBUTES (decl))) + { + error_at (loc, "variable %qD used more than once with " + "%<#pragma acc declare%>", decl); + error = true; + continue; + } + + if (!error) + { + tree id; + + if (OMP_CLAUSE_MAP_KIND (t) == GOMP_MAP_LINK) + id = get_identifier ("omp declare target link"); + else + id = get_identifier ("omp declare target"); + + DECL_ATTRIBUTES (decl) + = tree_cons (id, NULL_TREE, DECL_ATTRIBUTES (decl)); + + if (global_bindings_p ()) + { + symtab_node *node = symtab_node::get (decl); + if (node != NULL) + { + node->offloadable = 1; +#ifdef ENABLE_OFFLOADING + g->have_offload = true; + if (is_a (node)) + { + vec_safe_push (offload_vars, decl); + node->force_output = 1; + } +#endif + } + } + } + } + + if (error || global_bindings_p ()) + return; + + stmt = make_node (OACC_DECLARE); + TREE_TYPE (stmt) = void_type_node; + OACC_DECLARE_CLAUSES (stmt) = clauses; + SET_EXPR_LOCATION (stmt, pragma_loc); + + add_stmt (stmt); + + return; +} + +/* OpenACC 2.0: # pragma acc enter data oacc-enter-data-clause[optseq] new-line or diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c index a87675e..0ab5275 100644 --- a/gcc/cp/parser.c +++ b/gcc/cp/parser.c @@ -29128,6 +29128,8 @@ cp_parser_omp_clause_name (cp_parser *parser) result = PRAGMA_OMP_CLAUSE_DEVICE; else if (!strcmp ("deviceptr", p)) result = PRAGMA_OACC_CLAUSE_DEVICEPTR; + else if (!strcmp ("device_resident", p)) + result = PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT; else if (!strcmp ("dist_schedule", p)) result = PRAGMA_OMP_CLAUSE_DIST_SCHEDULE; break; @@ -29541,10 +29543,16 @@ cp_parser_oacc_data_clause (cp_parser *parser, pragma_omp_clause c_kind, case PRAGMA_OACC_CLAUSE_DEVICE: kind = GOMP_MAP_FORCE_TO; break; + case PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT: + kind = GOMP_MAP_DEVICE_RESIDENT; + break; case PRAGMA_OACC_CLAUSE_HOST: case PRAGMA_OACC_CLAUSE_SELF: kind = GOMP_MAP_FORCE_FROM; break; + case PRAGMA_OACC_CLAUSE_LINK: + kind = GOMP_MAP_LINK; + break; case PRAGMA_OACC_CLAUSE_PRESENT: kind = GOMP_MAP_FORCE_PRESENT; break; @@ -31545,6 +31553,10 @@ cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask, clauses = cp_parser_oacc_data_clause_deviceptr (parser, clauses); c_name = "deviceptr"; break; + case PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT: + clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses); + c_name = "device_resident"; + break; case PRAGMA_OACC_CLAUSE_FIRSTPRIVATE: clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE_FIRSTPRIVATE, clauses); @@ -31569,6 +31581,10 @@ cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask, clauses, here); c_name = "independent"; break; + case PRAGMA_OACC_CLAUSE_LINK: + clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses); + c_name = "link"; + break; case PRAGMA_OACC_CLAUSE_NUM_GANGS: code = OMP_CLAUSE_NUM_GANGS; c_name = "num_gangs"; @@ -34526,6 +34542,158 @@ cp_parser_oacc_data (cp_parser *parser, cp_token *pragma_tok) } /* OpenACC 2.0: + # pragma acc declare oacc-data-clause[optseq] new-line +*/ + +#define OACC_DECLARE_CLAUSE_MASK \ + ( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPY) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYIN) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYOUT) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_CREATE) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICEPTR) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_LINK) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE)) + +static tree +cp_parser_oacc_declare (cp_parser *parser, cp_token *pragma_tok) +{ + tree clauses, stmt, t; + bool error = false; + + clauses = cp_parser_oacc_all_clauses (parser, OACC_DECLARE_CLAUSE_MASK, + "#pragma acc declare", pragma_tok, true); + + + if (find_omp_clause (clauses, OMP_CLAUSE_MAP) == NULL_TREE) + { + error_at (pragma_tok->location, + "no valid clauses specified in %<#pragma acc declare%>"); + return NULL_TREE; + } + + for (tree t = clauses; t; t = OMP_CLAUSE_CHAIN (t)) + { + location_t loc = OMP_CLAUSE_LOCATION (t); + tree decl = OMP_CLAUSE_DECL (t); + if (!DECL_P (decl)) + { + error_at (loc, "array section in %<#pragma acc declare%>"); + error = true; + continue; + } + gcc_assert (OMP_CLAUSE_CODE (t) == OMP_CLAUSE_MAP); + switch (OMP_CLAUSE_MAP_KIND (t)) + { + case GOMP_MAP_FORCE_ALLOC: + case GOMP_MAP_FORCE_TO: + case GOMP_MAP_FORCE_DEVICEPTR: + case GOMP_MAP_DEVICE_RESIDENT: + break; + + case GOMP_MAP_POINTER: + /* Generated by c_finish_omp_clauses from array sections; + avoid spurious diagnostics. */ + break; + + case GOMP_MAP_LINK: + if (!global_bindings_p () + && (TREE_STATIC (decl) + || !DECL_EXTERNAL (decl))) + { + error_at (loc, + "%qD must be a global variable in" + "%<#pragma acc declare link%>", + decl); + error = true; + continue; + } + break; + + default: + if (global_bindings_p ()) + { + error_at (loc, "invalid OpenACC clause at file scope"); + error = true; + continue; + } + if (DECL_EXTERNAL (decl)) + { + error_at (loc, + "invalid use of % variable %qD " + "in %<#pragma acc declare%>", decl); + error = true; + continue; + } + else if (TREE_PUBLIC (decl)) + { + error_at (loc, + "invalid use of % variable %qD " + "in %<#pragma acc declare%>", decl); + error = true; + continue; + } + break; + } + + if (lookup_attribute ("omp declare target", DECL_ATTRIBUTES (decl)) + || lookup_attribute ("omp declare target link", + DECL_ATTRIBUTES (decl))) + { + error_at (loc, "variable %qD used more than once with " + "%<#pragma acc declare%>", decl); + error = true; + continue; + } + + if (!error) + { + tree id; + + if (OMP_CLAUSE_MAP_KIND (t) == GOMP_MAP_LINK) + id = get_identifier ("omp declare target link"); + else + id = get_identifier ("omp declare target"); + + DECL_ATTRIBUTES (decl) + = tree_cons (id, NULL_TREE, DECL_ATTRIBUTES (decl)); + if (global_bindings_p ()) + { + symtab_node *node = symtab_node::get (decl); + if (node != NULL) + { + node->offloadable = 1; +#ifdef ENABLE_OFFLOADING + g->have_offload = true; + if (is_a (node)) + { + vec_safe_push (offload_vars, decl); + node->force_output = 1; + } +#endif + } + } + } + } + + if (error || global_bindings_p ()) + return NULL_TREE; + + stmt = make_node (OACC_DECLARE); + TREE_TYPE (stmt) = void_type_node; + OACC_DECLARE_CLAUSES (stmt) = clauses; + SET_EXPR_LOCATION (stmt, pragma_tok->location); + + add_stmt (stmt); + + return NULL_TREE; +} + +/* OpenACC 2.0: # pragma acc enter data oacc-enter-data-clause[optseq] new-line or @@ -36354,6 +36522,10 @@ cp_parser_pragma (cp_parser *parser, enum pragma_context context) cp_parser_omp_declare (parser, pragma_tok, context); return false; + case PRAGMA_OACC_DECLARE: + cp_parser_oacc_declare (parser, pragma_tok); + return false; + case PRAGMA_OACC_ROUTINE: cp_parser_oacc_routine (parser, pragma_tok, context); return false; diff --git a/gcc/cp/pt.c b/gcc/cp/pt.c index 053a4ef..18d0561 100644 --- a/gcc/cp/pt.c +++ b/gcc/cp/pt.c @@ -15403,6 +15403,14 @@ tsubst_expr (tree t, tree args, tsubst_flags_t complain, tree in_decl, add_stmt (t); break; + case OACC_DECLARE: + t = copy_node (t); + tmp = tsubst_omp_clauses (OACC_DECLARE_CLAUSES (t), false, false, + args, complain, in_decl); + OACC_DECLARE_CLAUSES (t) = tmp; + add_stmt (t); + break; + case OMP_TARGET_UPDATE: case OMP_TARGET_ENTER_DATA: case OMP_TARGET_EXIT_DATA: diff --git a/gcc/gimple-pretty-print.c b/gcc/gimple-pretty-print.c index 7b50cdf..7764201 100644 --- a/gcc/gimple-pretty-print.c +++ b/gcc/gimple-pretty-print.c @@ -1353,6 +1353,9 @@ dump_gimple_omp_target (pretty_printer *buffer, gomp_target *gs, case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA: kind = " oacc_enter_exit_data"; break; + case GF_OMP_TARGET_KIND_OACC_DECLARE: + kind = " oacc_declare"; + break; default: gcc_unreachable (); } diff --git a/gcc/gimple.h b/gcc/gimple.h index 781801b..e45162d 100644 --- a/gcc/gimple.h +++ b/gcc/gimple.h @@ -170,6 +170,7 @@ enum gf_mask { GF_OMP_TARGET_KIND_OACC_DATA = 7, GF_OMP_TARGET_KIND_OACC_UPDATE = 8, GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA = 9, + GF_OMP_TARGET_KIND_OACC_DECLARE = 10, /* True on an GIMPLE_OMP_RETURN statement if the return does not require a thread synchronization via some sort of barrier. The exact barrier @@ -6004,6 +6005,7 @@ is_gimple_omp_oacc (const gimple *stmt) case GF_OMP_TARGET_KIND_OACC_DATA: case GF_OMP_TARGET_KIND_OACC_UPDATE: case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA: + case GF_OMP_TARGET_KIND_OACC_DECLARE: return true; default: return false; diff --git a/gcc/gimplify.c b/gcc/gimplify.c index 66e5168..7a7458d 100644 --- a/gcc/gimplify.c +++ b/gcc/gimplify.c @@ -176,6 +176,7 @@ static struct gimplify_omp_ctx *gimplify_omp_ctxp; /* Forward declaration. */ static enum gimplify_status gimplify_compound_expr (tree *, gimple_seq *, bool); +static hash_map *oacc_declare_returns; /* Shorter alias name for the above function for use in gimplify.c only. */ @@ -1078,6 +1079,7 @@ gimplify_bind_expr (tree *expr_p, gimple_seq *pre_p) gimple_seq body, cleanup; gcall *stack_save; location_t start_locus = 0, end_locus = 0; + tree ret_clauses = NULL; tree temp = voidify_wrapper_expr (bind_expr, NULL); @@ -1179,9 +1181,39 @@ gimplify_bind_expr (tree *expr_p, gimple_seq *pre_p) clobber_stmt = gimple_build_assign (t, clobber); gimple_set_location (clobber_stmt, end_locus); gimplify_seq_add_stmt (&cleanup, clobber_stmt); + + if (flag_openacc && oacc_declare_returns != NULL) + { + tree *c = oacc_declare_returns->get (t); + if (c != NULL) + { + if (ret_clauses) + OMP_CLAUSE_CHAIN (*c) = ret_clauses; + + ret_clauses = *c; + + oacc_declare_returns->remove (t); + + if (oacc_declare_returns->elements () == 0) + { + delete oacc_declare_returns; + oacc_declare_returns = NULL; + } + } + } } } + if (ret_clauses) + { + gomp_target *stmt; + gimple_stmt_iterator si = gsi_start (cleanup); + + stmt = gimple_build_omp_target (NULL, GF_OMP_TARGET_KIND_OACC_DECLARE, + ret_clauses); + gsi_insert_seq_before_without_update (&si, stmt, GSI_NEW_STMT); + } + if (cleanup) { gtry *gs; @@ -5809,6 +5841,26 @@ omp_notice_threadprivate_variable (struct gimplify_omp_ctx *ctx, tree decl, return false; } +/* Return true if global var DECL is device resident. */ + +static bool +device_resident_p (tree decl) +{ + tree attr = lookup_attribute ("oacc declare target", DECL_ATTRIBUTES (decl)); + + if (!attr) + return false; + + for (tree t = TREE_VALUE (attr); t; t = TREE_PURPOSE (t)) + { + tree c = TREE_VALUE (t); + if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DEVICE_RESIDENT) + return true; + } + + return false; +} + /* Determine outer default flags for DECL mentioned in an OMP region but not declared in an enclosing clause. @@ -5858,6 +5910,10 @@ omp_default_clause (struct gimplify_omp_ctx *ctx, tree decl, flags |= GOVD_FIRSTPRIVATE; break; case OMP_CLAUSE_DEFAULT_UNSPECIFIED: + if (is_global_var (decl) + && ctx->region_type & (ORT_ACC_PARALLEL | ORT_ACC_KERNELS) + && device_resident_p (decl)) + flags |= GOVD_MAP_TO_ONLY | GOVD_MAP; /* decl will be either GOVD_FIRSTPRIVATE or GOVD_SHARED. */ gcc_assert ((ctx->region_type & ORT_TASK) != 0); if (struct gimplify_omp_ctx *octx = ctx->outer_context) @@ -7763,6 +7819,121 @@ gimplify_oacc_cache (tree *expr_p, gimple_seq *pre_p) *expr_p = NULL_TREE; } +/* Helper function of gimplify_oacc_declare. The helper's purpose is to, + if required, translate 'kind' in CLAUSE into an 'entry' kind and 'exit' + kind. The entry kind will replace the one in CLAUSE, while the exit + kind will be used in a new omp_clause and returned to the caller. */ + +static tree +gimplify_oacc_declare_1 (tree clause) +{ + HOST_WIDE_INT kind, new_op; + bool ret = false; + tree c = NULL; + + kind = OMP_CLAUSE_MAP_KIND (clause); + + switch (kind) + { + case GOMP_MAP_ALLOC: + case GOMP_MAP_FORCE_ALLOC: + case GOMP_MAP_FORCE_TO: + new_op = GOMP_MAP_FORCE_DEALLOC; + 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; + ret = true; + break; + + case GOMP_MAP_FROM: + OMP_CLAUSE_SET_MAP_KIND (clause, GOMP_MAP_FORCE_ALLOC); + new_op = GOMP_MAP_FROM; + ret = true; + break; + + case GOMP_MAP_TOFROM: + OMP_CLAUSE_SET_MAP_KIND (clause, GOMP_MAP_TO); + new_op = GOMP_MAP_FROM; + ret = true; + break; + + case GOMP_MAP_DEVICE_RESIDENT: + case GOMP_MAP_FORCE_DEVICEPTR: + case GOMP_MAP_FORCE_PRESENT: + case GOMP_MAP_LINK: + case GOMP_MAP_POINTER: + case GOMP_MAP_TO: + break; + + default: + gcc_unreachable (); + break; + } + + if (ret) + { + c = build_omp_clause (OMP_CLAUSE_LOCATION (clause), OMP_CLAUSE_MAP); + OMP_CLAUSE_SET_MAP_KIND (c, new_op); + OMP_CLAUSE_DECL (c) = OMP_CLAUSE_DECL (clause); + } + + return c; +} + +/* Gimplify OACC_DECLARE. */ + +static void +gimplify_oacc_declare (tree *expr_p, gimple_seq *pre_p) +{ + tree expr = *expr_p; + gomp_target *stmt; + tree clauses, t; + + clauses = OACC_DECLARE_CLAUSES (expr); + + gimplify_scan_omp_clauses (&clauses, pre_p, ORT_TARGET_DATA, OACC_DECLARE); + + for (t = clauses; t; t = OMP_CLAUSE_CHAIN (t)) + { + tree decl = OMP_CLAUSE_DECL (t); + + if (TREE_CODE (decl) == MEM_REF) + continue; + + if (TREE_CODE (decl) == VAR_DECL + && !is_global_var (decl) + && DECL_CONTEXT (decl) == current_function_decl) + { + tree c = gimplify_oacc_declare_1 (t); + if (c) + { + if (oacc_declare_returns == NULL) + oacc_declare_returns = new hash_map; + + oacc_declare_returns->put (decl, c); + } + } + + omp_add_variable (gimplify_omp_ctxp, decl, GOVD_SEEN); + } + + stmt = gimple_build_omp_target (NULL, GF_OMP_TARGET_KIND_OACC_DECLARE, + clauses); + + gimplify_seq_add_stmt (pre_p, stmt); + + *expr_p = NULL_TREE; +} + /* Gimplify the contents of an OMP_PARALLEL statement. This involves gimplification of the body, as well as scanning the body for used variables. We need to do this scan now, because variable-sized @@ -10123,11 +10294,15 @@ gimplify_expr (tree *expr_p, gimple_seq *pre_p, gimple_seq *post_p, break; case OACC_HOST_DATA: - case OACC_DECLARE: sorry ("directive not yet implemented"); ret = GS_ALL_DONE; break; + case OACC_DECLARE: + gimplify_oacc_declare (expr_p, pre_p); + ret = GS_ALL_DONE; + break; + case OACC_DATA: case OACC_KERNELS: case OACC_PARALLEL: diff --git a/gcc/omp-builtins.def b/gcc/omp-builtins.def index 0b6bd58..d540dab 100644 --- a/gcc/omp-builtins.def +++ b/gcc/omp-builtins.def @@ -353,3 +353,5 @@ DEF_GOMP_BUILTIN (BUILT_IN_GOMP_TARGET_ENTER_EXIT_DATA, BT_FN_VOID_INT_SIZE_PTR_PTR_PTR_UINT_PTR, ATTR_NOTHROW_LIST) DEF_GOMP_BUILTIN (BUILT_IN_GOMP_TEAMS, "GOMP_teams", BT_FN_VOID_UINT_UINT, ATTR_NOTHROW_LIST) +DEF_GOACC_BUILTIN (BUILT_IN_GOACC_DECLARE, "GOACC_declare", + BT_FN_VOID_INT_SIZE_PTR_PTR_PTR, ATTR_NOTHROW_LIST) diff --git a/gcc/omp-low.c b/gcc/omp-low.c index 51b471c..f7584de 100644 --- a/gcc/omp-low.c +++ b/gcc/omp-low.c @@ -12454,6 +12454,7 @@ expand_omp_target (struct omp_region *region) case GF_OMP_TARGET_KIND_OACC_KERNELS: case GF_OMP_TARGET_KIND_OACC_UPDATE: case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA: + case GF_OMP_TARGET_KIND_OACC_DECLARE: data_region = false; break; case GF_OMP_TARGET_KIND_DATA: @@ -12697,6 +12698,9 @@ expand_omp_target (struct omp_region *region) case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA: start_ix = BUILT_IN_GOACC_ENTER_EXIT_DATA; break; + case GF_OMP_TARGET_KIND_OACC_DECLARE: + start_ix = BUILT_IN_GOACC_DECLARE; + break; default: gcc_unreachable (); } @@ -12819,6 +12823,7 @@ expand_omp_target (struct omp_region *region) switch (start_ix) { case BUILT_IN_GOACC_DATA_START: + case BUILT_IN_GOACC_DECLARE: case BUILT_IN_GOMP_TARGET_DATA: break; case BUILT_IN_GOMP_TARGET: @@ -13133,6 +13138,7 @@ build_omp_regions_1 (basic_block bb, struct omp_region *parent, case GF_OMP_TARGET_KIND_EXIT_DATA: case GF_OMP_TARGET_KIND_OACC_UPDATE: case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA: + case GF_OMP_TARGET_KIND_OACC_DECLARE: /* ..., other than for those stand-alone directives... */ region = NULL; break; @@ -14916,6 +14922,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) case GF_OMP_TARGET_KIND_OACC_KERNELS: case GF_OMP_TARGET_KIND_OACC_UPDATE: case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA: + case GF_OMP_TARGET_KIND_OACC_DECLARE: data_region = false; break; case GF_OMP_TARGET_KIND_DATA: @@ -14987,6 +14994,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) case GOMP_MAP_FORCE_TOFROM: case GOMP_MAP_FORCE_PRESENT: case GOMP_MAP_FORCE_DEVICEPTR: + case GOMP_MAP_DEVICE_RESIDENT: + case GOMP_MAP_LINK: gcc_assert (is_gimple_omp_oacc (stmt)); break; default: @@ -16713,6 +16722,7 @@ make_gimple_omp_edges (basic_block bb, struct omp_region **region, case GF_OMP_TARGET_KIND_EXIT_DATA: case GF_OMP_TARGET_KIND_OACC_UPDATE: case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA: + case GF_OMP_TARGET_KIND_OACC_DECLARE: cur_region = cur_region->outer; break; default: diff --git a/gcc/testsuite/c-c++-common/goacc/declare-1.c b/gcc/testsuite/c-c++-common/goacc/declare-1.c new file mode 100644 index 0000000..b036c63 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/declare-1.c @@ -0,0 +1,83 @@ +/* Test valid uses of declare directive. */ +/* { dg-do compile } */ + +int v0; +#pragma acc declare create(v0) + +int v1; +#pragma acc declare copyin(v1) + +int *v2; +#pragma acc declare deviceptr(v2) + +int v3; +#pragma acc declare device_resident(v3) + +int v4; +#pragma acc declare link(v4) + +int v5, v6, v7, v8; +#pragma acc declare create(v5, v6) copyin(v7, v8) + +void +f (void) +{ + int va0; +#pragma acc declare create(va0) + + int va1; +#pragma acc declare copyin(va1) + + int *va2; +#pragma acc declare deviceptr(va2) + + int va3; +#pragma acc declare device_resident(va3) + + extern int ve0; +#pragma acc declare create(ve0) + + extern int ve1; +#pragma acc declare copyin(ve1) + + extern int *ve2; +#pragma acc declare deviceptr(ve2) + + extern int ve3; +#pragma acc declare device_resident(ve3) + + extern int ve4; +#pragma acc declare link(ve4) + + int va5; +#pragma acc declare copy(va5) + + int va6; +#pragma acc declare copyout(va6) + + int va7; +#pragma acc declare present(va7) + + int va8; +#pragma acc declare present_or_copy(va8) + + int va9; +#pragma acc declare present_or_copyin(va9) + + int va10; +#pragma acc declare present_or_copyout(va10) + + int va11; +#pragma acc declare present_or_create(va11) + + a: + { + int va0; +#pragma acc declare create(va0) + if (v1) + goto a; + else + goto b; + } + b:; +} diff --git a/gcc/testsuite/c-c++-common/goacc/declare-2.c b/gcc/testsuite/c-c++-common/goacc/declare-2.c new file mode 100644 index 0000000..d24cb22 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/declare-2.c @@ -0,0 +1,79 @@ +/* Test invalid uses of declare directive. */ +/* { dg-do compile } */ + +#pragma acc declare /* { dg-error "no valid clauses" } */ + +#pragma acc declare create(undeclared) /* { dg-error "undeclared" } */ +/* { dg-error "no valid clauses" "second error" { target *-*-* } 6 } */ + +int v0[10]; +#pragma acc declare create(v0[1:3]) /* { dg-error "array section" } */ + +int v1; +#pragma acc declare create(v1, v1) /* { dg-error "more than once" } */ + +int v2; +#pragma acc declare create(v2) +#pragma acc declare copyin(v2) /* { dg-error "more than once" } */ + +int v3; +#pragma acc declare copy(v3) /* { dg-error "at file scope" } */ + +int v4; +#pragma acc declare copyout(v4) /* { dg-error "at file scope" } */ + +int v5; +#pragma acc declare present(v5) /* { dg-error "at file scope" } */ + +int v6; +#pragma acc declare present_or_copy(v6) /* { dg-error "at file scope" } */ + +int v7; +#pragma acc declare present_or_copyin(v7) /* { dg-error "at file scope" } */ + +int v8; +#pragma acc declare present_or_copyout(v8) /* { dg-error "at file scope" } */ + +int v9; +#pragma acc declare present_or_create(v9) /* { dg-error "at file scope" } */ + +int va10; +#pragma acc declare create (va10) +#pragma acc declare link (va10) /* { dg-error "more than once" } */ + +int va11; +#pragma acc declare link (va11) +#pragma acc declare link (va11) /* { dg-error "more than once" } */ + +int va12; +#pragma acc declare create (va12) link (va12) /* { dg-error "more than once" } */ + +void +f (void) +{ + int va0; +#pragma acc declare link(va0) /* { dg-error "global variable" } */ + + extern int ve0; +#pragma acc declare copy(ve0) /* { dg-error "invalid use of" } */ + + extern int ve1; +#pragma acc declare copyout(ve1) /* { dg-error "invalid use of" } */ + + extern int ve2; +#pragma acc declare present(ve2) /* { dg-error "invalid use of" } */ + + extern int ve3; +#pragma acc declare present_or_copy(ve3) /* { dg-error "invalid use of" } */ + + extern int ve4; +#pragma acc declare present_or_copyin(ve4) /* { dg-error "invalid use of" } */ + + extern int ve5; +#pragma acc declare present_or_copyout(ve5) /* { dg-error "invalid use of" } */ + + extern int ve6; +#pragma acc declare present_or_create(ve6) /* { dg-error "invalid use of" } */ + +#pragma acc declare present (v9) /* { dg-error "invalid use of" } */ +} diff --git a/gcc/tree-pretty-print.c b/gcc/tree-pretty-print.c index 3f0a4e6..caec760 100644 --- a/gcc/tree-pretty-print.c +++ b/gcc/tree-pretty-print.c @@ -654,6 +654,12 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, int flags) case GOMP_MAP_ALWAYS_POINTER: pp_string (pp, "always_pointer"); break; + case GOMP_MAP_DEVICE_RESIDENT: + pp_string (pp, "device_resident"); + break; + case GOMP_MAP_LINK: + pp_string (pp, "link"); + break; default: gcc_unreachable (); } diff --git a/include/gomp-constants.h b/include/gomp-constants.h index 7671dd7..dffd631 100644 --- a/include/gomp-constants.h +++ b/include/gomp-constants.h @@ -72,6 +72,11 @@ enum gomp_map_kind POINTER_SIZE_UNITS. */ GOMP_MAP_FORCE_DEVICEPTR = (GOMP_MAP_FLAG_SPECIAL_1 | 0), /* Do not map, copy bits for firstprivate instead. */ + /* OpenACC device_resident. */ + GOMP_MAP_DEVICE_RESIDENT = (GOMP_MAP_FLAG_SPECIAL_1 | 1), + /* OpenACC link. */ + GOMP_MAP_LINK = (GOMP_MAP_FLAG_SPECIAL_1 | 2), + /* Allocate. */ GOMP_MAP_FIRSTPRIVATE = (GOMP_MAP_FLAG_SPECIAL | 0), /* Similarly, but store the value in the pointer rather than pointed by the pointer. */ diff --git a/libgomp/libgomp.map b/libgomp/libgomp.map index 39faba9..d16710f 100644 --- a/libgomp/libgomp.map +++ b/libgomp/libgomp.map @@ -392,6 +392,7 @@ GOACC_2.0 { GOACC_2.0.1 { global: + GOACC_declare; GOACC_parallel_keyed; } GOACC_2.0; diff --git a/libgomp/oacc-parallel.c b/libgomp/oacc-parallel.c index 525846b..f76943a 100644 --- a/libgomp/oacc-parallel.c +++ b/libgomp/oacc-parallel.c @@ -501,3 +501,61 @@ GOACC_get_thread_num (void) { return 0; } + +void +GOACC_declare (int device, size_t mapnum, + void **hostaddrs, size_t *sizes, unsigned short *kinds) +{ + int i; + + for (i = 0; i < mapnum; i++) + { + unsigned char kind = kinds[i] & 0xff; + + if (kind == GOMP_MAP_POINTER || kind == GOMP_MAP_TO_PSET) + continue; + + switch (kind) + { + case GOMP_MAP_FORCE_ALLOC: + case GOMP_MAP_FORCE_DEALLOC: + case GOMP_MAP_FORCE_FROM: + case GOMP_MAP_FORCE_TO: + case GOMP_MAP_POINTER: + GOACC_enter_exit_data (device, 1, &hostaddrs[i], &sizes[i], + &kinds[i], 0, 0); + break; + + case GOMP_MAP_FORCE_DEVICEPTR: + break; + + case GOMP_MAP_ALLOC: + if (!acc_is_present (hostaddrs[i], sizes[i])) + GOACC_enter_exit_data (device, 1, &hostaddrs[i], &sizes[i], + &kinds[i], 0, 0); + break; + + case GOMP_MAP_TO: + GOACC_enter_exit_data (device, 1, &hostaddrs[i], &sizes[i], + &kinds[i], 0, 0); + + break; + + case GOMP_MAP_FROM: + kinds[i] = GOMP_MAP_FORCE_FROM; + GOACC_enter_exit_data (device, 1, &hostaddrs[i], &sizes[i], + &kinds[i], 0, 0); + break; + + case GOMP_MAP_FORCE_PRESENT: + if (!acc_is_present (hostaddrs[i], sizes[i])) + gomp_fatal ("[%p,%ld] is not mapped", hostaddrs[i], + (unsigned long) sizes[i]); + break; + + default: + assert (0); + break; + } + } +} diff --git a/libgomp/testsuite/libgomp.oacc-c++/declare-1.C b/libgomp/testsuite/libgomp.oacc-c++/declare-1.C new file mode 100644 index 0000000..0286955 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c++/declare-1.C @@ -0,0 +1,31 @@ +/* { dg-do run { target openacc_nvidia_accel_selected } } */ + +#include + +template +T foo() +{ + T a, b; + #pragma acc declare create (a) + + #pragma acc parallel copyout (b) + { + a = 5; + b = a; + } + + return b; +} + +int +main (void) +{ + int rc; + + rc = foo(); + + if (rc != 5) + abort (); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-1.c new file mode 100644 index 0000000..c63a68d --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-1.c @@ -0,0 +1,122 @@ +/* { dg-do run { target openacc_nvidia_accel_selected } } */ + +#include +#include +#include + +#define N 8 + +void +subr2 (int *a) +{ + int i; + int f[N]; +#pragma acc declare copyout (f) + +#pragma acc parallel copy (a[0:N]) + { + for (i = 0; i < N; i++) + { + f[i] = a[i]; + a[i] = f[i] + f[i] + f[i]; + } + } +} + +void +subr1 (int *a) +{ + int f[N]; +#pragma acc declare copy (f) + +#pragma acc parallel copy (a[0:N]) + { + int i; + + for (i = 0; i < N; i++) + { + f[i] = a[i]; + a[i] = f[i] + f[i]; + } + } +} + +int b[8]; +#pragma acc declare create (b) + +int d[8] = { 1, 2, 3, 4, 5, 6, 7, 8 }; +#pragma acc declare copyin (d) + +int +main (int argc, char **argv) +{ + int a[N]; + int e[N]; +#pragma acc declare create (e) + int i; + + for (i = 0; i < N; i++) + a[i] = i + 1; + + if (!acc_is_present (&b, sizeof (b))) + abort (); + + if (!acc_is_present (&d, sizeof (d))) + abort (); + + if (!acc_is_present (&e, sizeof (e))) + abort (); + +#pragma acc parallel copyin (a[0:N]) + { + for (i = 0; i < N; i++) + { + b[i] = a[i]; + a[i] = b[i]; + } + } + + for (i = 0; i < N; i++) + { + if (a[i] != i + 1) + abort (); + } + +#pragma acc parallel copy (a[0:N]) + { + for (i = 0; i < N; i++) + { + e[i] = a[i] + d[i]; + a[i] = e[i]; + } + } + + for (i = 0; i < N; i++) + { + if (a[i] != (i + 1) * 2) + abort (); + } + + for (i = 0; i < N; i++) + { + a[i] = 1234; + } + + subr1 (&a[0]); + + for (i = 0; i < N; i++) + { + if (a[i] != 1234 * 2) + abort (); + } + + subr2 (&a[0]); + + for (i = 0; i < 1; i++) + { + if (a[i] != 1234 * 6) + abort (); + } + + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-2.c new file mode 100644 index 0000000..2078a33 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-2.c @@ -0,0 +1,64 @@ +/* { dg-do run { target openacc_nvidia_accel_selected } } */ + +#include + +#define N 16 + +float c[N]; +#pragma acc declare device_resident (c) + +#pragma acc routine +float +subr2 (float a) +{ + int i; + + for (i = 0; i < N; i++) + c[i] = 2.0; + + for (i = 0; i < N; i++) + a += c[i]; + + return a; +} + +float b[N]; +#pragma acc declare copyin (b) + +#pragma acc routine +float +subr1 (float a) +{ + int i; + + for (i = 0; i < N; i++) + a += b[i]; + + return a; +} + +int +main (int argc, char **argv) +{ + float a; + int i; + + for (i = 0; i < 16; i++) + b[i] = 1.0; + + a = 0.0; + + a = subr1 (a); + + if (a != 16.0) + abort (); + + a = 0.0; + + a = subr2 (a); + + if (a != 32.0) + abort (); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-4.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-4.c new file mode 100644 index 0000000..013310e --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-4.c @@ -0,0 +1,41 @@ +/* { dg-do run { target openacc_nvidia_accel_selected } } */ + +#include +#include + +float b; +#pragma acc declare link (b) + +#pragma acc routine +int +func (int a) +{ + b = a + 1; + + return b; +} + +int +main (int argc, char **argv) +{ + float a; + + a = 2.0; + +#pragma acc parallel copy (a) + { + b = a; + a = 1.0; + a = a + b; + } + + if (a != 3.0) + abort (); + + a = func (a); + + if (a != 4.0) + abort (); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-5.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-5.c new file mode 100644 index 0000000..38c5de0 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-5.c @@ -0,0 +1,15 @@ +/* { dg-do run { target openacc_nvidia_accel_selected } } */ + +#include + +int +main (int argc, char **argv) +{ + int a[8] __attribute__((unused)); + + fprintf (stderr, "CheCKpOInT\n"); +#pragma acc declare present (a) +} + +/* { dg-output "CheCKpOInT" } */ +/* { dg-shouldfail "" } */