From patchwork Tue Jul 14 15:58:56 2015 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Jakub Jelinek X-Patchwork-Id: 495163 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 90BDB14076E for ; Wed, 15 Jul 2015 01:59:13 +1000 (AEST) Authentication-Results: ozlabs.org; dkim=pass (1024-bit key; unprotected) header.d=gcc.gnu.org header.i=@gcc.gnu.org header.b=TSgDraEF; 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:date :from:to:cc:subject:message-id:reply-to:mime-version :content-type; q=dns; s=default; b=uf3ck7saZjhDtZH8rhmCPjkmcsq83 1g8DVQU5e+sLOzEe73CqDS5o0UDwyWV2uCV47eCkowcX9atjzLPZ0+oi4oDmbLVm oVotiB9NzISR5mZDAdeWipYmvsa65HfzcjaJhWRlqueyRnCIkCjOP8jb5WaByil3 sYxAhuxplWwH9g= 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:date :from:to:cc:subject:message-id:reply-to:mime-version :content-type; s=default; bh=YlhN2UWJ46DnVluzHzTQRTlaFUE=; b=TSg DraEF0WCzG2yrGkqb936jnJrVC26+sjZDpMcT9Die3Pr4PfPMxvI8dZOZ4NAqUod 0OO+k1Cf+qGx2Yrxz7fSJPMaZzmoHGS+QBd0S9pzIFme0kXEnUsvCY1R7qtoXLBg EZ/iW6sv9S+AL+oZULsWFQ3h4WTOpRqZlpj1lARk= Received: (qmail 98448 invoked by alias); 14 Jul 2015 15:59:05 -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 98439 invoked by uid 89); 14 Jul 2015 15:59:04 -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, KAM_LAZY_DOMAIN_SECURITY, RP_MATCHES_RCVD, SPF_HELO_PASS autolearn=ham version=3.3.2 X-HELO: mx1.redhat.com Received: from mx1.redhat.com (HELO mx1.redhat.com) (209.132.183.28) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with (AES256-GCM-SHA384 encrypted) ESMTPS; Tue, 14 Jul 2015 15:59:02 +0000 Received: from int-mx10.intmail.prod.int.phx2.redhat.com (int-mx10.intmail.prod.int.phx2.redhat.com [10.5.11.23]) by mx1.redhat.com (Postfix) with ESMTPS id BE84F3187D7; Tue, 14 Jul 2015 15:59:00 +0000 (UTC) Received: from tucnak.zalov.cz (ovpn-116-44.ams2.redhat.com [10.36.116.44]) by int-mx10.intmail.prod.int.phx2.redhat.com (8.14.4/8.14.4) with ESMTP id t6EFwwx8021250 (version=TLSv1/SSLv3 cipher=DHE-RSA-AES256-GCM-SHA384 bits=256 verify=NO); Tue, 14 Jul 2015 11:58:59 -0400 Received: from tucnak.zalov.cz (localhost [127.0.0.1]) by tucnak.zalov.cz (8.14.9/8.14.9) with ESMTP id t6EFwv2e023819; Tue, 14 Jul 2015 17:58:57 +0200 Received: (from jakub@localhost) by tucnak.zalov.cz (8.14.9/8.14.9/Submit) id t6EFwu4H023818; Tue, 14 Jul 2015 17:58:56 +0200 Date: Tue, 14 Jul 2015 17:58:56 +0200 From: Jakub Jelinek To: Ilya Verbin Cc: gcc-patches@gcc.gnu.org Subject: [gomp4.1] Parsing of {use,is}_device_ptr clauses Message-ID: <20150714155856.GJ1788@tucnak.redhat.com> Reply-To: Jakub Jelinek MIME-Version: 1.0 Content-Disposition: inline User-Agent: Mutt/1.5.23 (2014-03-12) X-IsSubscribed: yes Hi! I've committed this patch to parse/gimplify/omp lower {use,is}_device_ptr clauses which I understood are meant to act like firstprivate clause with special side-effects, but they are ignored for now during OpenMP expansion. I'm committing this now, since it still shouldn't break the tests. I've skipped parsing of struct members in map/to/from clauses for now, because it is too unclear (#383) to me, waiting for clarifications. The next step, adjusting gimplifier to the new rules, will break some of the tests, and then we need to find out how to implement firstprivate clause on target construct, how to implement newly the pointer translation for array sections base (that should now be handled like private with special action), how to implement is_device_ptr (supposedly just as firstprivate - I think this clause got added just because of OpenCL) and how to implement use_device_ptr. 2015-07-14 Jakub Jelinek * tree-core.h (enum omp_clause_code): Add OMP_CLAUSE_USE_DEVICE_PTR and OMP_CLAUSE_IS_DEVICE_PTR. * gimplify.c (gimplify_scan_omp_clauses): Handle OMP_CLAUSE_USE_DEVICE_PTR and OMP_CLAUSE_IS_DEVICE_PTR like OMP_CLAUSE_FIRSTPRIVATE. (gimplify_adjust_omp_clauses): Handle OMP_CLAUSE_{USE,IS}_DEVICE_PTR. * tree.c (omp_clause_num_ops): Add entries for OMP_CLAUSE_{USE,IS}_DEVICE_PTR. (omp_clause_code_name): Likewise. (walk_tree_1): Handle OMP_CLAUSE_{USE,IS}_DEVICE_PTR. * tree-nested.c (convert_nonlocal_omp_clauses, convert_local_omp_clauses): Likewise. * omp-low.c (scan_sharing_clauses): Likewise. * tree-pretty-print.c (dump_omp_clause): Likewise. c-family/ * c-omp.c (c_omp_split_clauses): Improve function comment. Handle OMP_CLAUSE_IS_DEVICE_PTR. Adjust for OMP_CLAUSE_PRIVATE and OMP_CLAUSE_FIRSTPRIVATE being supported also on #pragma omp target construct. * c-pragma.h (enum pragma_omp_clause): Add PRAGMA_OMP_CLAUSE_IS_DEVICE_PTR and PRAGMA_OMP_CLAUSE_USE_DEVICE_PTR. c/ * c-parser.c (c_parser_omp_clause_name): Parse use_device_ptr and is_device_ptr. (c_parser_omp_clause_use_device_ptr, c_parser_omp_clause_is_device_ptr): New functions. (c_parser_omp_all_clauses): Handle PRAGMA_OMP_CLAUSE_{IS,USE}_DEVICE_PTR. (OMP_TARGET_DATA_CLAUSE_MASK): Add PRAGMA_OMP_CLAUSE_USE_DEVICE_PTR. (OMP_TARGET_CLAUSE_MASK): Add PRAGMA_OMP_CLAUSE_IS_DEVICE_PTR. * c-typeck.c (c_finish_omp_clauses): Handle OMP_CLAUSE_{USE,IS}_DEVICE_PTR. cp/ * parser.c (cp_parser_omp_clause_name): Parse use_device_ptr and is_device_ptr. (cp_parser_omp_all_clauses): Handle OMP_CLAUSE_{USE,IS}_DEVICE_PTR. (OMP_TARGET_DATA_CLAUSE_MASK): Add PRAGMA_OMP_CLAUSE_USE_DEVICE_PTR. (OMP_TARGET_CLAUSE_MASK): Add PRAGMA_OMP_CLAUSE_IS_DEVICE_PTR. * pt.c (tsubst_omp_clauses): Handle OMP_CLAUSE_{USE,IS}_DEVICE_PTR. * semantics.c (finish_omp_clauses): Likewise. Jakub --- gcc/gimplify.c.jj 2015-07-14 14:29:49.274536887 +0200 +++ gcc/gimplify.c 2015-07-14 16:17:52.993914085 +0200 @@ -6394,6 +6394,13 @@ gimplify_scan_omp_clauses (tree *list_p, } goto do_notice; + case OMP_CLAUSE_USE_DEVICE_PTR: + flags = GOVD_FIRSTPRIVATE | GOVD_EXPLICIT; + goto do_add; + case OMP_CLAUSE_IS_DEVICE_PTR: + flags = GOVD_FIRSTPRIVATE | GOVD_EXPLICIT; + goto do_add; + do_add: decl = OMP_CLAUSE_DECL (c); do_add_decl: @@ -6957,6 +6964,8 @@ gimplify_adjust_omp_clauses (gimple_seq case OMP_CLAUSE_SIMD: case OMP_CLAUSE_HINT: case OMP_CLAUSE_DEFAULTMAP: + case OMP_CLAUSE_USE_DEVICE_PTR: + case OMP_CLAUSE_IS_DEVICE_PTR: case OMP_CLAUSE__CILK_FOR_COUNT_: case OMP_CLAUSE_ASYNC: case OMP_CLAUSE_WAIT: --- gcc/tree.c.jj 2015-07-14 14:29:49.032547791 +0200 +++ gcc/tree.c 2015-07-14 14:49:57.666673189 +0200 @@ -291,6 +291,8 @@ unsigned const char omp_clause_num_ops[] 2, /* OMP_CLAUSE_FROM */ 2, /* OMP_CLAUSE_TO */ 2, /* OMP_CLAUSE_MAP */ + 1, /* OMP_CLAUSE_USE_DEVICE_PTR */ + 1, /* OMP_CLAUSE_IS_DEVICE_PTR */ 2, /* OMP_CLAUSE__CACHE_ */ 1, /* OMP_CLAUSE_DEVICE_RESIDENT */ 1, /* OMP_CLAUSE_USE_DEVICE */ @@ -358,6 +360,8 @@ const char * const omp_clause_code_name[ "from", "to", "map", + "use_device_ptr", + "is_device_ptr", "_cache_", "device_resident", "use_device", @@ -11388,6 +11392,8 @@ walk_tree_1 (tree *tp, walk_tree_fn func case OMP_CLAUSE_GRAINSIZE: case OMP_CLAUSE_NUM_TASKS: case OMP_CLAUSE_HINT: + case OMP_CLAUSE_USE_DEVICE_PTR: + case OMP_CLAUSE_IS_DEVICE_PTR: case OMP_CLAUSE__LOOPTEMP_: case OMP_CLAUSE__SIMDUID_: case OMP_CLAUSE__CILK_FOR_COUNT_: --- gcc/tree-nested.c.jj 2015-07-14 14:29:49.329534408 +0200 +++ gcc/tree-nested.c 2015-07-14 14:49:57.667673145 +0200 @@ -1098,6 +1098,8 @@ convert_nonlocal_omp_clauses (tree *pcla case OMP_CLAUSE_FIRSTPRIVATE: case OMP_CLAUSE_COPYPRIVATE: case OMP_CLAUSE_SHARED: + case OMP_CLAUSE_USE_DEVICE_PTR: + case OMP_CLAUSE_IS_DEVICE_PTR: do_decl_clause: decl = OMP_CLAUSE_DECL (clause); if (TREE_CODE (decl) == VAR_DECL @@ -1743,6 +1745,8 @@ convert_local_omp_clauses (tree *pclause case OMP_CLAUSE_FIRSTPRIVATE: case OMP_CLAUSE_COPYPRIVATE: case OMP_CLAUSE_SHARED: + case OMP_CLAUSE_USE_DEVICE_PTR: + case OMP_CLAUSE_IS_DEVICE_PTR: do_decl_clause: decl = OMP_CLAUSE_DECL (clause); if (TREE_CODE (decl) == VAR_DECL --- gcc/tree-core.h.jj 2015-07-14 14:29:49.061546484 +0200 +++ gcc/tree-core.h 2015-07-14 14:49:57.675672794 +0200 @@ -270,6 +270,12 @@ enum omp_clause_code { OpenMP clause: map ({alloc:,to:,from:,tofrom:,}variable-list). */ OMP_CLAUSE_MAP, + /* OpenMP clause: use_device_ptr (variable-list). */ + OMP_CLAUSE_USE_DEVICE_PTR, + + /* OpenMP clause: is_device_ptr (variable-list). */ + OMP_CLAUSE_IS_DEVICE_PTR, + /* Internal structure to hold OpenACC cache directive's variable-list. #pragma acc cache (variable-list). */ OMP_CLAUSE__CACHE_, --- gcc/omp-low.c.jj 2015-07-14 14:30:06.000000000 +0200 +++ gcc/omp-low.c 2015-07-14 16:27:54.944892284 +0200 @@ -1954,6 +1954,11 @@ scan_sharing_clauses (tree clauses, omp_ } break; + case OMP_CLAUSE_USE_DEVICE_PTR: + case OMP_CLAUSE_IS_DEVICE_PTR: + decl = OMP_CLAUSE_DECL (c); + goto do_private; + case OMP_CLAUSE__LOOPTEMP_: gcc_assert (is_taskreg_ctx (ctx)); decl = OMP_CLAUSE_DECL (c); @@ -2138,6 +2143,8 @@ scan_sharing_clauses (tree clauses, omp_ /* FALLTHRU */ case OMP_CLAUSE_PRIVATE: case OMP_CLAUSE_LINEAR: + case OMP_CLAUSE_USE_DEVICE_PTR: + case OMP_CLAUSE_IS_DEVICE_PTR: decl = OMP_CLAUSE_DECL (c); if (is_variable_sized (decl)) install_var_local (decl, ctx); --- gcc/tree-pretty-print.c.jj 2015-07-14 14:29:49.022548242 +0200 +++ gcc/tree-pretty-print.c 2015-07-14 14:49:57.676672750 +0200 @@ -329,6 +329,12 @@ dump_omp_clause (pretty_printer *pp, tre case OMP_CLAUSE_UNIFORM: name = "uniform"; goto print_remap; + case OMP_CLAUSE_USE_DEVICE_PTR: + name = "use_device_ptr"; + goto print_remap; + case OMP_CLAUSE_IS_DEVICE_PTR: + name = "is_device_ptr"; + goto print_remap; case OMP_CLAUSE__LOOPTEMP_: name = "_looptemp_"; goto print_remap; --- gcc/c-family/c-omp.c.jj 2015-07-14 14:29:49.281536571 +0200 +++ gcc/c-family/c-omp.c 2015-07-14 14:49:57.674672838 +0200 @@ -684,7 +684,7 @@ c_finish_omp_for (location_t locus, enum } } -/* Right now we have 15 different combined constructs, this +/* Right now we have 15 different combined/composite constructs, this function attempts to split or duplicate clauses for combined constructs. CODE is the innermost construct in the combined construct, and MASK allows to determine which constructs are combined together, @@ -744,6 +744,7 @@ c_omp_split_clauses (location_t loc, enu /* First the clauses that are unique to some constructs. */ case OMP_CLAUSE_DEVICE: case OMP_CLAUSE_MAP: + case OMP_CLAUSE_IS_DEVICE_PTR: s = C_OMP_CLAUSE_SPLIT_TARGET; break; case OMP_CLAUSE_NUM_TEAMS: @@ -814,7 +815,7 @@ c_omp_split_clauses (location_t loc, enu else s = C_OMP_CLAUSE_SPLIT_DISTRIBUTE; break; - /* Private clause is supported on all constructs but target, + /* Private clause is supported on all constructs, it is enough to put it on the innermost one. For #pragma omp {for,sections} put it on parallel though, as that's what we did for OpenMP 3.1. */ @@ -830,9 +831,18 @@ c_omp_split_clauses (location_t loc, enu } break; /* Firstprivate clause is supported on all constructs but - target and simd. Put it on the outermost of those and - duplicate on parallel. */ + simd. Put it on the outermost of those and duplicate on teams + and parallel. */ case OMP_CLAUSE_FIRSTPRIVATE: + if ((mask & (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_MAP)) + != 0) + { + c = build_omp_clause (OMP_CLAUSE_LOCATION (clauses), + OMP_CLAUSE_FIRSTPRIVATE); + OMP_CLAUSE_DECL (c) = OMP_CLAUSE_DECL (clauses); + OMP_CLAUSE_CHAIN (c) = cclauses[C_OMP_CLAUSE_SPLIT_TARGET]; + cclauses[C_OMP_CLAUSE_SPLIT_TARGET] = c; + } if ((mask & (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_NUM_THREADS)) != 0) { --- gcc/c-family/c-pragma.h.jj 2015-07-14 14:29:49.286536346 +0200 +++ gcc/c-family/c-pragma.h 2015-07-14 14:49:57.674672838 +0200 @@ -98,6 +98,7 @@ typedef enum pragma_omp_clause { PRAGMA_OMP_CLAUSE_HINT, PRAGMA_OMP_CLAUSE_IF, PRAGMA_OMP_CLAUSE_INBRANCH, + PRAGMA_OMP_CLAUSE_IS_DEVICE_PTR, PRAGMA_OMP_CLAUSE_LASTPRIVATE, PRAGMA_OMP_CLAUSE_LINEAR, PRAGMA_OMP_CLAUSE_MAP, @@ -126,6 +127,7 @@ typedef enum pragma_omp_clause { PRAGMA_OMP_CLAUSE_TO, PRAGMA_OMP_CLAUSE_UNIFORM, PRAGMA_OMP_CLAUSE_UNTIED, + PRAGMA_OMP_CLAUSE_USE_DEVICE_PTR, /* Clauses for Cilk Plus SIMD-enabled function. */ PRAGMA_CILK_CLAUSE_NOMASK, --- gcc/c/c-parser.c.jj 2015-07-14 14:29:49.338534003 +0200 +++ gcc/c/c-parser.c 2015-07-14 16:03:49.750735756 +0200 @@ -9945,6 +9945,8 @@ c_parser_omp_clause_name (c_parser *pars case 'i': if (!strcmp ("inbranch", p)) result = PRAGMA_OMP_CLAUSE_INBRANCH; + else if (!strcmp ("is_device_ptr", p)) + result = PRAGMA_OMP_CLAUSE_IS_DEVICE_PTR; break; case 'l': if (!strcmp ("lastprivate", p)) @@ -10045,6 +10047,8 @@ c_parser_omp_clause_name (c_parser *pars result = PRAGMA_OMP_CLAUSE_UNIFORM; else if (!strcmp ("untied", p)) result = PRAGMA_OMP_CLAUSE_UNTIED; + else if (!strcmp ("use_device_ptr", p)) + result = PRAGMA_OMP_CLAUSE_USE_DEVICE_PTR; break; case 'v': if (!strcmp ("vector", p)) @@ -10916,6 +10920,25 @@ c_parser_omp_clause_defaultmap (c_parser return list; } +/* OpenMP 4.1: + use_device_ptr ( variable-list ) */ + +static tree +c_parser_omp_clause_use_device_ptr (c_parser *parser, tree list) +{ + return c_parser_omp_var_list_parens (parser, OMP_CLAUSE_USE_DEVICE_PTR, + list); +} + +/* OpenMP 4.1: + is_device_ptr ( variable-list ) */ + +static tree +c_parser_omp_clause_is_device_ptr (c_parser *parser, tree list) +{ + return c_parser_omp_var_list_parens (parser, OMP_CLAUSE_IS_DEVICE_PTR, list); +} + /* OpenACC: num_workers ( expression ) */ @@ -12403,6 +12426,14 @@ c_parser_omp_all_clauses (c_parser *pars clauses = c_parser_omp_clause_map (parser, clauses); c_name = "map"; break; + case PRAGMA_OMP_CLAUSE_USE_DEVICE_PTR: + clauses = c_parser_omp_clause_use_device_ptr (parser, clauses); + c_name = "use_device_ptr"; + break; + case PRAGMA_OMP_CLAUSE_IS_DEVICE_PTR: + clauses = c_parser_omp_clause_is_device_ptr (parser, clauses); + c_name = "is_device_ptr"; + break; case PRAGMA_OMP_CLAUSE_DEVICE: clauses = c_parser_omp_clause_device (parser, clauses); c_name = "device"; @@ -14382,7 +14413,8 @@ c_parser_omp_teams (location_t loc, c_pa #define OMP_TARGET_DATA_CLAUSE_MASK \ ( (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEVICE) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_MAP) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IF)) + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IF) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_USE_DEVICE_PTR)) static tree c_parser_omp_target_data (location_t loc, c_parser *parser) @@ -14664,7 +14696,8 @@ c_parser_omp_target_exit_data (location_ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_NOWAIT) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PRIVATE) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_FIRSTPRIVATE) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEFAULTMAP)) + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEFAULTMAP) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IS_DEVICE_PTR)) static bool c_parser_omp_target (c_parser *parser, enum pragma_context context) --- gcc/c/c-typeck.c.jj 2015-07-14 14:29:49.349533507 +0200 +++ gcc/c/c-typeck.c 2015-07-14 16:55:50.130770229 +0200 @@ -12586,6 +12586,18 @@ c_finish_omp_clauses (tree clauses, bool } goto check_dup_generic; + case OMP_CLAUSE_IS_DEVICE_PTR: + case OMP_CLAUSE_USE_DEVICE_PTR: + t = OMP_CLAUSE_DECL (c); + if (TREE_CODE (TREE_TYPE (t)) != POINTER_TYPE) + { + error_at (OMP_CLAUSE_LOCATION (c), + "%qs variable is not a pointer", + omp_clause_code_name[OMP_CLAUSE_CODE (c)]); + remove = true; + } + goto check_dup_generic; + case OMP_CLAUSE_NOWAIT: if (copyprivate_seen) { --- gcc/cp/parser.c.jj 2015-07-14 14:29:49.311535219 +0200 +++ gcc/cp/parser.c 2015-07-14 17:39:06.281183207 +0200 @@ -27740,6 +27740,8 @@ cp_parser_omp_clause_name (cp_parser *pa case 'i': if (!strcmp ("inbranch", p)) result = PRAGMA_OMP_CLAUSE_INBRANCH; + else if (!strcmp ("is_device_ptr", p)) + result = PRAGMA_OMP_CLAUSE_IS_DEVICE_PTR; break; case 'l': if (!strcmp ("lastprivate", p)) @@ -27836,6 +27838,8 @@ cp_parser_omp_clause_name (cp_parser *pa result = PRAGMA_OMP_CLAUSE_UNIFORM; else if (!strcmp ("untied", p)) result = PRAGMA_OMP_CLAUSE_UNTIED; + else if (!strcmp ("use_device_ptr", p)) + result = PRAGMA_OMP_CLAUSE_USE_DEVICE_PTR; break; case 'v': if (!strcmp ("vector_length", p)) @@ -29852,6 +29856,16 @@ cp_parser_omp_all_clauses (cp_parser *pa token->location); c_name = "defaultmap"; break; + case PRAGMA_OMP_CLAUSE_USE_DEVICE_PTR: + clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE_USE_DEVICE_PTR, + clauses); + c_name = "use_device_ptr"; + break; + case PRAGMA_OMP_CLAUSE_IS_DEVICE_PTR: + clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE_IS_DEVICE_PTR, + clauses); + c_name = "is_device_ptr"; + break; case PRAGMA_OMP_CLAUSE_IF: clauses = cp_parser_omp_clause_if (parser, clauses, token->location); c_name = "if"; @@ -32008,7 +32022,8 @@ cp_parser_omp_teams (cp_parser *parser, #define OMP_TARGET_DATA_CLAUSE_MASK \ ( (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEVICE) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_MAP) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IF)) + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IF) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_USE_DEVICE_PTR)) static tree cp_parser_omp_target_data (cp_parser *parser, cp_token *pragma_tok) @@ -32292,7 +32307,8 @@ cp_parser_omp_target_update (cp_parser * | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_NOWAIT) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PRIVATE) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_FIRSTPRIVATE) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEFAULTMAP)) + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEFAULTMAP) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IS_DEVICE_PTR)) static bool cp_parser_omp_target (cp_parser *parser, cp_token *pragma_tok, --- gcc/cp/pt.c.jj 2015-07-14 14:29:49.299535760 +0200 +++ gcc/cp/pt.c 2015-07-14 14:49:57.659673496 +0200 @@ -13601,6 +13601,8 @@ tsubst_omp_clauses (tree clauses, bool d case OMP_CLAUSE_FROM: case OMP_CLAUSE_TO: case OMP_CLAUSE_MAP: + case OMP_CLAUSE_USE_DEVICE_PTR: + case OMP_CLAUSE_IS_DEVICE_PTR: OMP_CLAUSE_DECL (nc) = tsubst_omp_clause_decl (OMP_CLAUSE_DECL (oc), args, complain, in_decl); @@ -13688,6 +13690,8 @@ tsubst_omp_clauses (tree clauses, bool d case OMP_CLAUSE_COPYPRIVATE: case OMP_CLAUSE_LINEAR: case OMP_CLAUSE_REDUCTION: + case OMP_CLAUSE_USE_DEVICE_PTR: + case OMP_CLAUSE_IS_DEVICE_PTR: /* tsubst_expr on SCOPE_REF results in returning finish_non_static_data_member result. Undo that here. */ if (TREE_CODE (OMP_CLAUSE_DECL (oc)) == SCOPE_REF --- gcc/cp/semantics.c.jj 2015-07-14 14:29:49.324534634 +0200 +++ gcc/cp/semantics.c 2015-07-14 14:49:57.661673409 +0200 @@ -6385,6 +6385,26 @@ finish_omp_clauses (tree clauses, bool a } break; + case OMP_CLAUSE_IS_DEVICE_PTR: + case OMP_CLAUSE_USE_DEVICE_PTR: + field_ok = allow_fields; + t = OMP_CLAUSE_DECL (c); + if (!type_dependent_expression_p (t)) + { + tree type = TREE_TYPE (t); + if (TREE_CODE (type) != POINTER_TYPE + && (TREE_CODE (type) != REFERENCE_TYPE + || TREE_CODE (TREE_TYPE (type)) != POINTER_TYPE)) + { + error_at (OMP_CLAUSE_LOCATION (c), + "%qs variable is not a pointer or reference " + "to pointer", + omp_clause_code_name[OMP_CLAUSE_CODE (c)]); + remove = true; + } + } + goto check_dup_generic; + case OMP_CLAUSE_NOWAIT: case OMP_CLAUSE_ORDERED: case OMP_CLAUSE_DEFAULT: