From patchwork Mon Nov 9 16:01:32 2015 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: James Norris X-Patchwork-Id: 541817 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 C416B14017E for ; Tue, 10 Nov 2015 03:01:55 +1100 (AEDT) Authentication-Results: ozlabs.org; dkim=pass (1024-bit key; unprotected) header.d=gcc.gnu.org header.i=@gcc.gnu.org header.b=VT75XsQP; 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=JJEA/T9zR+4HNjusa CURzYR114CdgrBPmjIz4XKWieUSND+bWEVCHQaORGf5J+OvA6fFpw6feB2Oae7wi 8/vImSKOajlv2h6ZixtCbdTtB5a2cx+4HQBaytze/AhYkLiOBbdi53BcCdZMXkDx DBLddXy0L77KCqfQvGyO5iPg90= 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=R5+X3g80z8WzTURf9XYTNdz OaOk=; b=VT75XsQPkuNt4t75rHmKzXlZhN0URsLVpCwpiYzZoz77OBciWPKW5my wS/9Op3h9ir01T32f1Y5gISN+HHqPV5q4kqLbZsV4Z3EAPICQjMQTE+JXwvqITRi 0+2BzK1NXfknV+BcE72ZUV7dM1PpMCQrK/Nh3Kd9wdIYr8yN8p7s= Received: (qmail 67225 invoked by alias); 9 Nov 2015 16:01:43 -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 67207 invoked by uid 89); 9 Nov 2015 16:01:42 -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; Mon, 09 Nov 2015 16:01:38 +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 1Zvosw-00016u-TN from James_Norris@mentor.com ; Mon, 09 Nov 2015 08:01:34 -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; Mon, 9 Nov 2015 08:01:33 -0800 Message-ID: <5640C35C.9030907@codesourcery.com> Date: Mon, 9 Nov 2015 10:01:32 -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: James Norris , Jakub Jelinek 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> In-Reply-To: <563F6BA5.9020606@codesourcery.com> X-TagToolbar-Keys: D20151109100132171 Jakub, This is an update of the patch from: https://gcc.gnu.org/ml/gcc-patches/2015-11/msg00893.html. There was an unused variable in c/c-parser.c that was removed. I've included the ChangeLog as a convenience, but nothing was changed in the file. 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 (gimplify_bind_expr): Prepend 'exit' stmt to cleanup. * 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. 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-5.c: Likewise. diff --git a/gcc/c-family/c-pragma.c b/gcc/c-family/c-pragma.c index ac11838..cd0cc27 100644 --- a/gcc/c-family/c-pragma.c +++ b/gcc/c-family/c-pragma.c @@ -1207,6 +1207,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 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, @@ -151,6 +152,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, @@ -175,7 +177,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 23d0107..8edf745 100644 --- a/gcc/c/c-parser.c +++ b/gcc/c/c-parser.c @@ -1231,6 +1231,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 *); @@ -9697,6 +9698,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; @@ -9982,6 +9987,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; @@ -10418,10 +10425,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; @@ -12703,6 +12716,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"; @@ -12725,6 +12742,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"; @@ -13182,6 +13203,247 @@ 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 c, clauses, ret_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))) + { + error_at (loc, "variable %qD used more than once with " + "%<#pragma acc declare%>", decl); + error = true; + continue; + } + + if (!error) + { + tree id; + + if (OMP_CLAUSE_CODE (t) == OMP_CLAUSE_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; + + ret_clauses = NULL_TREE; + + for (c = clauses; c; c = OMP_CLAUSE_CHAIN (c)) + { + bool ret = false; + HOST_WIDE_INT kind, new_op; + + kind = OMP_CLAUSE_MAP_KIND (c); + + 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 (c, GOMP_MAP_FORCE_ALLOC); + new_op = GOMP_MAP_FORCE_FROM; + ret = true; + break; + + case GOMP_MAP_FORCE_TOFROM: + OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_FORCE_TO); + new_op = GOMP_MAP_FORCE_FROM; + ret = true; + break; + + case GOMP_MAP_FROM: + OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_FORCE_ALLOC); + new_op = GOMP_MAP_FROM; + ret = true; + break; + + case GOMP_MAP_TOFROM: + OMP_CLAUSE_SET_MAP_KIND (c, 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) + { + t = build_omp_clause (OMP_CLAUSE_LOCATION (c) , OMP_CLAUSE_MAP); + OMP_CLAUSE_SET_MAP_KIND (t, new_op); + OMP_CLAUSE_DECL (t) = OMP_CLAUSE_DECL (c); + + if (ret_clauses) + OMP_CLAUSE_CHAIN (t) = ret_clauses; + + ret_clauses = t; + } + } + + if (ret_clauses) + { + tree fndecl = current_function_decl; + tree attrs = lookup_attribute ("oacc declare returns", + DECL_ATTRIBUTES (fndecl)); + + if (attrs) + { + OMP_CLAUSE_CHAIN (ret_clauses) = TREE_VALUE (attrs); + TREE_VALUE (attrs) = ret_clauses; + } + else + { + tree id = get_identifier ("oacc declare returns"); + DECL_ATTRIBUTES (fndecl) = + tree_cons (id, ret_clauses, DECL_ATTRIBUTES (fndecl)); + + } + } + + 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 c6f5729..6432a34 100644 --- a/gcc/cp/parser.c +++ b/gcc/cp/parser.c @@ -29099,6 +29099,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; @@ -29512,10 +29514,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; @@ -31516,6 +31524,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); @@ -31540,6 +31552,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"; @@ -34497,6 +34513,246 @@ 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 c, clauses, ret_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))) + { + error_at (loc, "variable %qD used more than once with " + "%<#pragma acc declare%>", decl); + error = true; + continue; + } + + if (!error) + { + tree id; + + if (OMP_CLAUSE_CODE (t) == OMP_CLAUSE_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; + + ret_clauses = NULL_TREE; + + for (c = clauses; c; c = OMP_CLAUSE_CHAIN (c)) + { + bool ret = false; + HOST_WIDE_INT kind, new_op; + + kind = OMP_CLAUSE_MAP_KIND (c); + + 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 (c, GOMP_MAP_FORCE_ALLOC); + new_op = GOMP_MAP_FORCE_FROM; + ret = true; + break; + + case GOMP_MAP_FORCE_TOFROM: + OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_FORCE_TO); + new_op = GOMP_MAP_FORCE_FROM; + ret = true; + break; + + case GOMP_MAP_FROM: + OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_FORCE_ALLOC); + new_op = GOMP_MAP_FROM; + ret = true; + break; + + case GOMP_MAP_TOFROM: + OMP_CLAUSE_SET_MAP_KIND (c, 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_POINTER: + case GOMP_MAP_TO: + break; + + case GOMP_MAP_LINK: + continue; + + default: + gcc_unreachable (); + break; + } + + if (ret) + { + t = build_omp_clause (OMP_CLAUSE_LOCATION (c) , OMP_CLAUSE_MAP); + OMP_CLAUSE_SET_MAP_KIND (t, new_op); + OMP_CLAUSE_DECL (t) = OMP_CLAUSE_DECL (c); + + if (ret_clauses) + OMP_CLAUSE_CHAIN (t) = ret_clauses; + + ret_clauses = t; + } + } + + if (ret_clauses) + { + tree fndecl = current_function_decl; + tree attrs = lookup_attribute ("oacc declare returns", + DECL_ATTRIBUTES (fndecl)); + + if (attrs) + { + OMP_CLAUSE_CHAIN (ret_clauses) = TREE_VALUE (attrs); + TREE_VALUE (attrs) = ret_clauses; + } + else + { + tree id = get_identifier ("oacc declare returns"); + DECL_ATTRIBUTES (fndecl) = + tree_cons (id, ret_clauses, DECL_ATTRIBUTES (fndecl)); + + } + } + + 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 @@ -36183,6 +36439,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_ATOMIC: case PRAGMA_OACC_CACHE: case PRAGMA_OACC_DATA: diff --git a/gcc/cp/pt.c b/gcc/cp/pt.c index 45eda3a..3e03f02 100644 --- a/gcc/cp/pt.c +++ b/gcc/cp/pt.c @@ -15422,6 +15422,17 @@ 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; + tmp = tsubst_omp_clauses (OACC_DECLARE_RETURN_CLAUSES (t), false, false, + args, complain, in_decl); + OACC_DECLARE_RETURN_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 fa34858..a25f07c 100644 --- a/gcc/gimplify.c +++ b/gcc/gimplify.c @@ -1065,6 +1065,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); @@ -1166,9 +1167,56 @@ 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) + { + tree attrs = lookup_attribute ("oacc declare returns", + DECL_ATTRIBUTES (current_function_decl)); + tree clauses, c, c_next = NULL, c_prev = NULL; + + if (!attrs) + break; + + clauses = TREE_VALUE (attrs); + + for (c = clauses; c; c_prev = c, c = c_next) + { + c_next = OMP_CLAUSE_CHAIN (c); + + if (t == OMP_CLAUSE_DECL (c)) + { + if (ret_clauses) + OMP_CLAUSE_CHAIN (c) = ret_clauses; + + ret_clauses = c; + + if (c_prev == NULL) + clauses = c_next; + else + OMP_CLAUSE_CHAIN (c_prev) = c_next; + } + } + + if (clauses == NULL) + { + DECL_ATTRIBUTES (current_function_decl) = + remove_attribute ("oacc declare returns", + DECL_ATTRIBUTES (current_function_decl)); + } + } } } + 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; @@ -5792,6 +5840,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. @@ -5841,6 +5909,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; /* 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) @@ -7712,6 +7782,37 @@ gimplify_oacc_cache (tree *expr_p, gimple_seq *pre_p) *expr_p = NULL_TREE; } +/* 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; + + 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 @@ -10063,11 +10164,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 fc87a3f..0365bc4 100644 --- a/gcc/omp-builtins.def +++ b/gcc/omp-builtins.def @@ -357,3 +357,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 5ffb276..0119e44 100644 --- a/gcc/omp-low.c +++ b/gcc/omp-low.c @@ -12344,6 +12344,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: @@ -12587,6 +12588,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 (); } @@ -12709,6 +12713,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: @@ -13023,6 +13028,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; @@ -14806,6 +14812,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: @@ -14876,6 +14883,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: @@ -16542,6 +16551,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..7979f0c --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/declare-2.c @@ -0,0 +1,68 @@ +/* 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 "subarray" } */ + +int v1; +#pragma acc declare create(v1, v1) /* { dg-error "more than once" } */ + +int v2; +#pragma acc declare create(v2) /* { dg-message "previous directive" } */ +#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" } */ + +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.h b/gcc/tree.h index 6768b3b..a84d11a 100644 --- a/gcc/tree.h +++ b/gcc/tree.h @@ -1232,6 +1232,8 @@ extern void protected_set_expr_location (tree, location_t); #define OACC_DECLARE_CLAUSES(NODE) \ TREE_OPERAND (OACC_DECLARE_CHECK (NODE), 0) +#define OACC_DECLARE_RETURN_CLAUSES(NODE) \ + TREE_OPERAND (OACC_DECLARE_CHECK (NODE), 1) #define OACC_ENTER_DATA_CLAUSES(NODE) \ TREE_OPERAND (OACC_ENTER_DATA_CHECK (NODE), 0) 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..9de9e55 100644 --- a/libgomp/oacc-parallel.c +++ b/libgomp/oacc-parallel.c @@ -501,3 +501,62 @@ 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,%zd] is not mapped", hostaddrs[i], sizes[i]); + break; + + default: + assert (0); + break; + } + } +} diff --git a/libgomp/testsuite/declare-1.c b/libgomp/testsuite/declare-1.c new file mode 100644 index 0000000..8fbec4d --- /dev/null +++ b/libgomp/testsuite/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/declare-5.c b/libgomp/testsuite/declare-5.c new file mode 100644 index 0000000..1e2f6ce --- /dev/null +++ b/libgomp/testsuite/declare-5.c @@ -0,0 +1,13 @@ +/* { dg-do run { target openacc_nvidia_accel_selected } } */ + +int +main (int argc, char **argv) +{ + int a[8] __attribute__((unused)); + + __builtin_printf ("CheCKpOInT\n"); +#pragma acc declare present (a) +} + +/* { dg-output "CheCKpOInT" } */ +/* { dg-shouldfail "" } */