From patchwork Fri Jun 6 05:54:41 2014 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Thomas Schwinge X-Patchwork-Id: 356688 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 37448140088 for ; Fri, 6 Jun 2014 15:55:18 +1000 (EST) DomainKey-Signature: a=rsa-sha1; c=nofws; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender:from :to:cc:subject:date:message-id:mime-version:content-type; q=dns; s=default; b=da1C/j3wxTPuZegy7C84gN6mPfp6DRTth2DNa7BVYK/dy3PUXx pH4zNIUd5UBf2E5P6UN9RQ4wIkCBFH+YDpcw5WLPvSxNHWNL7SBJeoiKo3DAJAtS E2onHnXPjOVPFsVReZCLF/FHzCNDEctOGP/jO8eQTZA9j6/o1J6Xvoivk= DKIM-Signature: v=1; a=rsa-sha1; c=relaxed; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender:from :to:cc:subject:date:message-id:mime-version:content-type; s= default; bh=LxjqboNZ02rHZseByaHHvP/3WGk=; b=YpN9Iwql+lFW7N9MLmyv JkN1XMm9Ioaqd8ADRzORhCUgGLyHzWZrjoGdCV7uWCHqzC4Tg8W7zAuHQtTg81Db /r+h7/45axwXujTyKq3r9Onv4QoF9Z4nV5PjLGTcOucsIK1uXc1N1fLvd/9JEVY3 ONkQwXpL2G0EECKVxEfh0uY= Received: (qmail 23680 invoked by alias); 6 Jun 2014 05:55: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 23660 invoked by uid 89); 6 Jun 2014 05:55:08 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-2.0 required=5.0 tests=AWL, BAYES_00 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; Fri, 06 Jun 2014 05:55:06 +0000 Received: from svr-orw-exc-10.mgc.mentorg.com ([147.34.98.58]) by relay1.mentorg.com with esmtp id 1Wsn7F-0004jc-QK from Thomas_Schwinge@mentor.com for gcc-patches@gcc.gnu.org; Thu, 05 Jun 2014 22:55:01 -0700 Received: from SVR-IES-FEM-01.mgc.mentorg.com ([137.202.0.104]) by SVR-ORW-EXC-10.mgc.mentorg.com with Microsoft SMTPSVC(6.0.3790.4675); Thu, 5 Jun 2014 22:55:01 -0700 Received: from feldtkeller.schwinge.homeip.net (137.202.0.76) by SVR-IES-FEM-01.mgc.mentorg.com (137.202.0.104) with Microsoft SMTP Server id 14.2.247.3; Fri, 6 Jun 2014 06:55:00 +0100 From: Thomas Schwinge To: CC: Subject: [GOMP4, COMMITTED] OpenACC update directive. Date: Fri, 6 Jun 2014 07:54:41 +0200 Message-ID: <1402034081-14718-1-git-send-email-thomas@codesourcery.com> MIME-Version: 1.0 From: tschwinge gcc/c-family/ * c-pragma.c (oacc_pragmas): Add "update". * c-pragma.h (enum pragma_kind): Add PRAGMA_OACC_UPDATE. (enum pragma_omp_clause): Add PRAGMA_OMP_CLAUSE_HOST and PRAGMA_OMP_CLAUSE_SELF. gcc/c/ * c-parser.c (c_parser_pragma): Handle PRAGMA_OACC_UPDATE. (c_parser_omp_clause_name, c_parser_oacc_data_clause) (c_parser_oacc_all_clauses): Handle PRAGMA_OMP_CLAUSE_DEVICE, PRAGMA_OMP_CLAUSE_HOST, PRAGMA_OMP_CLAUSE_SELF. (c_parser_oacc_update): New function. gcc/ * gimple.h (enum gf_mask): Add GF_OMP_TARGET_KIND_OACC_UPDATE, and extend mask size, GF_OMP_TARGET_KIND_MASK. (is_gimple_omp_oacc_specifically): Handle GF_OMP_TARGET_KIND_OACC_UPDATE. * gimplify.c (gimplify_omp_target_update, gimplify_expr): Likewise. * gimple-pretty-print.c (dump_gimple_omp_target): Likewise. * omp-low.c (scan_omp_target, expand_omp_target) (build_omp_regions_1, lower_omp_target, lower_omp_1) (make_gimple_omp_edges): Likewise. * oacc-builtins.def (BUILT_IN_GOACC_UPDATE): New builtin. gcc/testsuite/ * c-c++-common/goacc/pragma_context.c: New file. * c-c++-common/goacc/update-1.c: Likewise. libgomp/ * libgomp.map (GOACC_2.0): Add GOACC_update. * oacc-parallel.c (GOACC_update): New function. git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gomp-4_0-branch@211299 138bc75d-0d04-0410-961f-82ee72b054a4 --- gcc/ChangeLog.gomp | 15 +++++ gcc/c-family/ChangeLog.gomp | 8 +++ gcc/c-family/c-pragma.c | 1 + gcc/c-family/c-pragma.h | 3 + gcc/c/ChangeLog.gomp | 9 +++ gcc/c/c-parser.c | 80 +++++++++++++++++++++++ gcc/gimple-pretty-print.c | 3 + gcc/gimple.h | 5 +- gcc/gimplify.c | 29 +++++--- gcc/oacc-builtins.def | 2 + gcc/omp-low.c | 30 ++++++--- gcc/testsuite/ChangeLog.gomp | 5 ++ gcc/testsuite/c-c++-common/goacc/pragma_context.c | 32 +++++++++ gcc/testsuite/c-c++-common/goacc/update-1.c | 10 +++ libgomp/ChangeLog.gomp | 6 ++ libgomp/libgomp.map | 1 + libgomp/oacc-parallel.c | 24 +++++++ 17 files changed, 246 insertions(+), 17 deletions(-) create mode 100644 gcc/testsuite/c-c++-common/goacc/pragma_context.c create mode 100644 gcc/testsuite/c-c++-common/goacc/update-1.c diff --git gcc/ChangeLog.gomp gcc/ChangeLog.gomp index 88f09b3..be1aa16 100644 --- gcc/ChangeLog.gomp +++ gcc/ChangeLog.gomp @@ -1,3 +1,18 @@ +2014-06-06 Thomas Schwinge + James Norris + + * gimple.h (enum gf_mask): Add GF_OMP_TARGET_KIND_OACC_UPDATE, and + extend mask size, GF_OMP_TARGET_KIND_MASK. + (is_gimple_omp_oacc_specifically): Handle + GF_OMP_TARGET_KIND_OACC_UPDATE. + * gimplify.c (gimplify_omp_target_update, gimplify_expr): + Likewise. + * gimple-pretty-print.c (dump_gimple_omp_target): Likewise. + * omp-low.c (scan_omp_target, expand_omp_target) + (build_omp_regions_1, lower_omp_target, lower_omp_1) + (make_gimple_omp_edges): Likewise. + * oacc-builtins.def (BUILT_IN_GOACC_UPDATE): New builtin. + 2014-06-05 Thomas Schwinge * gimplify.c (gimplify_scan_omp_clauses) diff --git gcc/c-family/ChangeLog.gomp gcc/c-family/ChangeLog.gomp index 37ebfe9..0894675 100644 --- gcc/c-family/ChangeLog.gomp +++ gcc/c-family/ChangeLog.gomp @@ -1,3 +1,11 @@ +2014-06-06 Thomas Schwinge + James Norris + + * c-pragma.c (oacc_pragmas): Add "update". + * c-pragma.h (enum pragma_kind): Add PRAGMA_OACC_UPDATE. + (enum pragma_omp_clause): Add PRAGMA_OMP_CLAUSE_HOST and + PRAGMA_OMP_CLAUSE_SELF. + 2014-03-20 Thomas Schwinge * c-omp.c (check_omp_for_incr_expr, c_finish_omp_for): Update diff --git gcc/c-family/c-pragma.c gcc/c-family/c-pragma.c index 85da7ac..6a9d8e0 100644 --- gcc/c-family/c-pragma.c +++ gcc/c-family/c-pragma.c @@ -1177,6 +1177,7 @@ static const struct omp_pragma_def oacc_pragmas[] = { { "kernels", PRAGMA_OACC_KERNELS }, { "loop", PRAGMA_OACC_LOOP }, { "parallel", PRAGMA_OACC_PARALLEL }, + { "update", PRAGMA_OACC_UPDATE }, }; static const struct omp_pragma_def omp_pragmas[] = { { "atomic", PRAGMA_OMP_ATOMIC }, diff --git gcc/c-family/c-pragma.h gcc/c-family/c-pragma.h index bb9c367..f3dfc63 100644 --- gcc/c-family/c-pragma.h +++ gcc/c-family/c-pragma.h @@ -31,6 +31,7 @@ typedef enum pragma_kind { PRAGMA_OACC_KERNELS, PRAGMA_OACC_LOOP, PRAGMA_OACC_PARALLEL, + PRAGMA_OACC_UPDATE, PRAGMA_OMP_ATOMIC, PRAGMA_OMP_BARRIER, PRAGMA_OMP_CANCEL, @@ -88,6 +89,7 @@ typedef enum pragma_omp_clause { PRAGMA_OMP_CLAUSE_FIRSTPRIVATE, PRAGMA_OMP_CLAUSE_FOR, PRAGMA_OMP_CLAUSE_FROM, + PRAGMA_OMP_CLAUSE_HOST, PRAGMA_OMP_CLAUSE_IF, PRAGMA_OMP_CLAUSE_INBRANCH, PRAGMA_OMP_CLAUSE_LASTPRIVATE, @@ -113,6 +115,7 @@ typedef enum pragma_omp_clause { PRAGMA_OMP_CLAUSE_SAFELEN, PRAGMA_OMP_CLAUSE_SCHEDULE, PRAGMA_OMP_CLAUSE_SECTIONS, + PRAGMA_OMP_CLAUSE_SELF, PRAGMA_OMP_CLAUSE_SHARED, PRAGMA_OMP_CLAUSE_SIMDLEN, PRAGMA_OMP_CLAUSE_TASKGROUP, diff --git gcc/c/ChangeLog.gomp gcc/c/ChangeLog.gomp index 1e80031..f1e45f3 100644 --- gcc/c/ChangeLog.gomp +++ gcc/c/ChangeLog.gomp @@ -1,3 +1,12 @@ +2014-06-06 Thomas Schwinge + James Norris + + * c-parser.c (c_parser_pragma): Handle PRAGMA_OACC_UPDATE. + (c_parser_omp_clause_name, c_parser_oacc_data_clause) + (c_parser_oacc_all_clauses): Handle PRAGMA_OMP_CLAUSE_DEVICE, + PRAGMA_OMP_CLAUSE_HOST, PRAGMA_OMP_CLAUSE_SELF. + (c_parser_oacc_update): New function. + 2014-06-05 Thomas Schwinge * c-typeck.c (handle_omp_array_sections, c_finish_omp_clauses): diff --git gcc/c/c-parser.c gcc/c/c-parser.c index e20348e..bf4bad62 100644 --- gcc/c/c-parser.c +++ gcc/c/c-parser.c @@ -1207,6 +1207,7 @@ static vec *c_parser_expr_list (c_parser *, bool, bool, static tree c_parser_oacc_loop (location_t, c_parser *, char *); static void c_parser_omp_construct (c_parser *); static void c_parser_omp_threadprivate (c_parser *); +static void c_parser_oacc_update (c_parser *); static void c_parser_omp_barrier (c_parser *); static void c_parser_omp_flush (c_parser *); static tree c_parser_omp_for_loop (location_t, c_parser *, enum tree_code, @@ -4457,6 +4458,14 @@ c_parser_initval (c_parser *parser, struct c_expr *after, Although they are erroneous if the labels declared aren't defined, is it useful for the syntax to be this way? + OpenACC: + + block-item: + openacc-directive + + openacc-directive: + update-directive + OpenMP: block-item: @@ -9433,6 +9442,17 @@ c_parser_pragma (c_parser *parser, enum pragma_context context) switch (id) { + case PRAGMA_OACC_UPDATE: + if (context != pragma_compound) + { + if (context == pragma_stmt) + c_parser_error (parser, "%<#pragma acc update%> may only be " + "used in compound statements"); + goto bad_stmt; + } + c_parser_oacc_update (parser); + return false; + case PRAGMA_OMP_BARRIER: if (context != pragma_compound) { @@ -9680,6 +9700,10 @@ c_parser_omp_clause_name (c_parser *parser) else if (!strcmp ("from", p)) result = PRAGMA_OMP_CLAUSE_FROM; break; + case 'h': + if (!strcmp ("host", p)) + result = PRAGMA_OMP_CLAUSE_SELF; + break; case 'i': if (!strcmp ("inbranch", p)) result = PRAGMA_OMP_CLAUSE_INBRANCH; @@ -9755,6 +9779,8 @@ c_parser_omp_clause_name (c_parser *parser) result = PRAGMA_OMP_CLAUSE_SHARED; else if (!strcmp ("simdlen", p)) result = PRAGMA_OMP_CLAUSE_SIMDLEN; + else if (!strcmp ("self", p)) + result = PRAGMA_OMP_CLAUSE_SELF; break; case 't': if (!strcmp ("taskgroup", p)) @@ -9960,6 +9986,12 @@ c_parser_oacc_data_clause (c_parser *parser, pragma_omp_clause c_kind, case PRAGMA_OMP_CLAUSE_DELETE: kind = OMP_CLAUSE_MAP_FORCE_DEALLOC; break; + case PRAGMA_OMP_CLAUSE_DEVICE: + kind = OMP_CLAUSE_MAP_FORCE_TO; + break; + case PRAGMA_OMP_CLAUSE_HOST: + kind = OMP_CLAUSE_MAP_FORCE_FROM; + break; case PRAGMA_OMP_CLAUSE_PRESENT: kind = OMP_CLAUSE_MAP_FORCE_PRESENT; break; @@ -9975,6 +10007,9 @@ c_parser_oacc_data_clause (c_parser *parser, pragma_omp_clause c_kind, case PRAGMA_OMP_CLAUSE_PRESENT_OR_CREATE: kind = OMP_CLAUSE_MAP_ALLOC; break; + case PRAGMA_OMP_CLAUSE_SELF: + kind = OMP_CLAUSE_MAP_FORCE_FROM; + break; } tree nl, c; nl = c_parser_omp_var_list_parens (parser, OMP_CLAUSE_MAP, list); @@ -11248,10 +11283,18 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask, clauses = c_parser_oacc_data_clause (parser, c_kind, clauses); c_name = "delete"; break; + case PRAGMA_OMP_CLAUSE_DEVICE: + clauses = c_parser_oacc_data_clause (parser, c_kind, clauses); + c_name = "device"; + break; case PRAGMA_OMP_CLAUSE_DEVICEPTR: clauses = c_parser_oacc_data_clause_deviceptr (parser, clauses); c_name = "deviceptr"; break; + case PRAGMA_OMP_CLAUSE_HOST: + clauses = c_parser_oacc_data_clause (parser, c_kind, clauses); + c_name = "host"; + break; case PRAGMA_OMP_CLAUSE_NUM_GANGS: clauses = c_parser_omp_clause_num_gangs (parser, clauses); c_name = "num_gangs"; @@ -11280,6 +11323,10 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask, clauses = c_parser_oacc_data_clause (parser, c_kind, clauses); c_name = "present_or_create"; break; + case PRAGMA_OMP_CLAUSE_SELF: + clauses = c_parser_oacc_data_clause (parser, c_kind, clauses); + c_name = "self"; + break; case PRAGMA_OMP_CLAUSE_VECTOR_LENGTH: clauses = c_parser_omp_clause_vector_length (parser, clauses); c_name = "vector_length"; @@ -11721,6 +11768,39 @@ c_parser_oacc_parallel (location_t loc, c_parser *parser, char *p_name) return stmt; } +/* OpenACC 2.0: + # pragma acc update oacc-update-clause[optseq] new-line +*/ + +#define OACC_UPDATE_CLAUSE_MASK \ + ( (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEVICE) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_HOST) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_SELF) ) + +static void +c_parser_oacc_update (c_parser *parser) +{ + location_t loc = c_parser_peek_token (parser)->location; + + c_parser_consume_pragma (parser); + + tree clauses = c_parser_oacc_all_clauses (parser, OACC_UPDATE_CLAUSE_MASK, + "#pragma acc update"); + if (find_omp_clause (clauses, OMP_CLAUSE_MAP) == NULL_TREE) + { + error_at (loc, + "%<#pragma acc update%> must contain at least one " + "% or % clause"); + return; + } + + tree stmt = make_node (OACC_UPDATE); + TREE_TYPE (stmt) = void_type_node; + OACC_UPDATE_CLAUSES (stmt) = clauses; + SET_EXPR_LOCATION (stmt, loc); + add_stmt (stmt); +} + /* OpenMP 2.5: # pragma omp atomic new-line expression-stmt diff --git gcc/gimple-pretty-print.c gcc/gimple-pretty-print.c index 53cb138..ab06551 100644 --- gcc/gimple-pretty-print.c +++ gcc/gimple-pretty-print.c @@ -1299,6 +1299,9 @@ dump_gimple_omp_target (pretty_printer *buffer, gimple gs, int spc, int flags) case GF_OMP_TARGET_KIND_OACC_DATA: kind = " oacc_data"; break; + case GF_OMP_TARGET_KIND_OACC_UPDATE: + kind = " oacc_update"; + break; default: gcc_unreachable (); } diff --git gcc/gimple.h gcc/gimple.h index 60b4896..1a1b952 100644 --- gcc/gimple.h +++ gcc/gimple.h @@ -101,11 +101,12 @@ enum gf_mask { GF_OMP_FOR_KIND_CILKSIMD = GF_OMP_FOR_SIMD | 1, GF_OMP_FOR_COMBINED = 1 << 3, GF_OMP_FOR_COMBINED_INTO = 1 << 4, - GF_OMP_TARGET_KIND_MASK = (1 << 2) - 1, + GF_OMP_TARGET_KIND_MASK = (1 << 3) - 1, GF_OMP_TARGET_KIND_REGION = 0, GF_OMP_TARGET_KIND_DATA = 1, GF_OMP_TARGET_KIND_UPDATE = 2, GF_OMP_TARGET_KIND_OACC_DATA = 3, + GF_OMP_TARGET_KIND_OACC_UPDATE = 4, /* True on an GIMPLE_OMP_RETURN statement if the return does not require a thread synchronization via some sort of barrier. The exact barrier @@ -5830,6 +5831,8 @@ is_gimple_omp_oacc_specifically (const_gimple stmt) { case GF_OMP_TARGET_KIND_OACC_DATA: return true; + case GF_OMP_TARGET_KIND_OACC_UPDATE: + return true; default: return false; } diff --git gcc/gimplify.c gcc/gimplify.c index a1b6be6..6b8e376 100644 --- gcc/gimplify.c +++ gcc/gimplify.c @@ -7229,19 +7229,32 @@ gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p) *expr_p = NULL_TREE; } -/* Gimplify the gross structure of OpenMP target update construct. */ +/* Gimplify the gross structure of OpenACC update and OpenMP target update + constructs. */ static void gimplify_omp_target_update (tree *expr_p, gimple_seq *pre_p) { - tree expr = *expr_p; + tree expr = *expr_p, clauses; + int kind; gimple stmt; - gimplify_scan_omp_clauses (&OMP_TARGET_UPDATE_CLAUSES (expr), pre_p, - ORT_WORKSHARE); - gimplify_adjust_omp_clauses (&OMP_TARGET_UPDATE_CLAUSES (expr)); - stmt = gimple_build_omp_target (NULL, GF_OMP_TARGET_KIND_UPDATE, - OMP_TARGET_UPDATE_CLAUSES (expr)); + switch (TREE_CODE (expr)) + { + case OACC_UPDATE: + clauses = OACC_UPDATE_CLAUSES (expr); + kind = GF_OMP_TARGET_KIND_OACC_UPDATE; + break; + case OMP_TARGET_UPDATE: + clauses = OMP_TARGET_UPDATE_CLAUSES (expr); + kind = GF_OMP_TARGET_KIND_UPDATE; + break; + default: + gcc_unreachable (); + } + gimplify_scan_omp_clauses (&clauses, pre_p, ORT_WORKSHARE); + gimplify_adjust_omp_clauses (&clauses); + stmt = gimple_build_omp_target (NULL, kind, clauses); gimplify_seq_add_stmt (pre_p, stmt); *expr_p = NULL_TREE; @@ -8169,7 +8182,6 @@ gimplify_expr (tree *expr_p, gimple_seq *pre_p, gimple_seq *post_p, case OACC_HOST_DATA: case OACC_DECLARE: - case OACC_UPDATE: case OACC_ENTER_DATA: case OACC_EXIT_DATA: case OACC_WAIT: @@ -8222,6 +8234,7 @@ gimplify_expr (tree *expr_p, gimple_seq *pre_p, gimple_seq *post_p, ret = GS_ALL_DONE; break; + case OACC_UPDATE: case OMP_TARGET_UPDATE: gimplify_omp_target_update (expr_p, pre_p); ret = GS_ALL_DONE; diff --git gcc/oacc-builtins.def gcc/oacc-builtins.def index 3935da4..dfb688c 100644 --- gcc/oacc-builtins.def +++ gcc/oacc-builtins.def @@ -37,3 +37,5 @@ DEF_GOACC_BUILTIN (BUILT_IN_GOACC_KERNELS, "GOACC_kernels", DEF_GOACC_BUILTIN (BUILT_IN_GOACC_PARALLEL, "GOACC_parallel", BT_FN_VOID_INT_OMPFN_PTR_SIZE_PTR_PTR_PTR_INT_INT_INT, ATTR_NOTHROW_LIST) +DEF_GOACC_BUILTIN (BUILT_IN_GOACC_UPDATE, "GOACC_update", + BT_FN_VOID_INT_PTR_SIZE_PTR_PTR_PTR, ATTR_NOTHROW_LIST) diff --git gcc/omp-low.c gcc/omp-low.c index 39f0598..6fd0f30 100644 --- gcc/omp-low.c +++ gcc/omp-low.c @@ -2419,7 +2419,8 @@ scan_omp_target (gimple stmt, omp_context *outer_ctx) tree name; int kind = gimple_omp_target_kind (stmt); - if (kind == GF_OMP_TARGET_KIND_OACC_DATA) + if (kind == GF_OMP_TARGET_KIND_OACC_DATA + || kind == GF_OMP_TARGET_KIND_OACC_UPDATE) { gcc_assert (taskreg_nesting_level == 0); gcc_assert (target_nesting_level == 0); @@ -8647,6 +8648,9 @@ expand_omp_target (struct omp_region *region) case GF_OMP_TARGET_KIND_OACC_DATA: start_ix = BUILT_IN_GOACC_DATA_START; break; + case GF_OMP_TARGET_KIND_OACC_UPDATE: + start_ix = BUILT_IN_GOACC_UPDATE; + break; default: gcc_unreachable (); } @@ -8662,9 +8666,11 @@ expand_omp_target (struct omp_region *region) cond = OMP_CLAUSE_IF_EXPR (c); c = find_omp_clause (clauses, OMP_CLAUSE_DEVICE); - gcc_assert (!c || kind != GF_OMP_TARGET_KIND_OACC_DATA); if (c) { + gcc_assert (kind != GF_OMP_TARGET_KIND_OACC_DATA + && kind != GF_OMP_TARGET_KIND_OACC_UPDATE); + device = OMP_CLAUSE_DEVICE_ID (c); clause_loc = OMP_CLAUSE_LOCATION (c); } @@ -8920,7 +8926,9 @@ build_omp_regions_1 (basic_block bb, struct omp_region *parent, ; } else if (code == GIMPLE_OMP_TARGET - && gimple_omp_target_kind (stmt) == GF_OMP_TARGET_KIND_UPDATE) + && (gimple_omp_target_kind (stmt) == GF_OMP_TARGET_KIND_UPDATE + || (gimple_omp_target_kind (stmt) + == GF_OMP_TARGET_KIND_OACC_UPDATE))) new_omp_region (bb, code, parent); else { @@ -10604,7 +10612,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) case OMP_CLAUSE_MAP_FORCE_PRESENT: case OMP_CLAUSE_MAP_FORCE_DEALLOC: case OMP_CLAUSE_MAP_FORCE_DEVICEPTR: - gcc_assert (kind == GF_OMP_TARGET_KIND_OACC_DATA); + gcc_assert (kind == GF_OMP_TARGET_KIND_OACC_DATA + || kind == GF_OMP_TARGET_KIND_OACC_UPDATE); break; default: gcc_unreachable (); @@ -10615,7 +10624,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) case OMP_CLAUSE_TO: case OMP_CLAUSE_FROM: if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP) - gcc_assert (kind != GF_OMP_TARGET_KIND_OACC_DATA); + gcc_assert (kind != GF_OMP_TARGET_KIND_OACC_DATA + && kind != GF_OMP_TARGET_KIND_OACC_UPDATE); var = OMP_CLAUSE_DECL (c); if (!DECL_P (var)) { @@ -10702,6 +10712,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) talign_shift = 3; break; case GF_OMP_TARGET_KIND_OACC_DATA: + case GF_OMP_TARGET_KIND_OACC_UPDATE: tkind_type = short_unsigned_type_node; talign_shift = 8; break; @@ -10904,7 +10915,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) else if (kind == GF_OMP_TARGET_KIND_DATA || kind == GF_OMP_TARGET_KIND_OACC_DATA) new_body = tgt_body; - if (kind != GF_OMP_TARGET_KIND_UPDATE) + if (kind != GF_OMP_TARGET_KIND_UPDATE + && kind != GF_OMP_TARGET_KIND_OACC_UPDATE) { gimple_seq_add_stmt (&new_body, gimple_build_omp_return (false)); gimple_omp_set_body (stmt, new_body); @@ -11124,7 +11136,8 @@ lower_omp_1 (gimple_stmt_iterator *gsi_p, omp_context *ctx) case GIMPLE_OMP_TARGET: ctx = maybe_lookup_ctx (stmt); gcc_assert (ctx); - if (gimple_omp_target_kind (stmt) == GF_OMP_TARGET_KIND_OACC_DATA) + if (gimple_omp_target_kind (stmt) == GF_OMP_TARGET_KIND_OACC_DATA + || gimple_omp_target_kind (stmt) == GF_OMP_TARGET_KIND_OACC_UPDATE) gcc_assert (!ctx->cancellable); lower_omp_target (gsi_p, ctx); break; @@ -11584,7 +11597,8 @@ make_gimple_omp_edges (basic_block bb, struct omp_region **region, case GIMPLE_OMP_TARGET: cur_region = new_omp_region (bb, code, cur_region); fallthru = true; - if (gimple_omp_target_kind (last) == GF_OMP_TARGET_KIND_UPDATE) + if (gimple_omp_target_kind (last) == GF_OMP_TARGET_KIND_UPDATE + || gimple_omp_target_kind (last) == GF_OMP_TARGET_KIND_OACC_UPDATE) cur_region = cur_region->outer; break; diff --git gcc/testsuite/ChangeLog.gomp gcc/testsuite/ChangeLog.gomp index 08ec907..17493ce 100644 --- gcc/testsuite/ChangeLog.gomp +++ gcc/testsuite/ChangeLog.gomp @@ -1,3 +1,8 @@ +2014-06-06 Thomas Schwinge + + * c-c++-common/goacc/pragma_context.c: New file. + * c-c++-common/goacc/update-1.c: Likewise. + 2014-06-05 Thomas Schwinge * c-c++-common/goacc/data-clause-duplicate-1.c: The OpenACC diff --git gcc/testsuite/c-c++-common/goacc/pragma_context.c gcc/testsuite/c-c++-common/goacc/pragma_context.c new file mode 100644 index 0000000..ad33d92 --- /dev/null +++ gcc/testsuite/c-c++-common/goacc/pragma_context.c @@ -0,0 +1,32 @@ +// pragma_external +#pragma acc update /* { dg-error "expected declaration specifiers before '#pragma'" } */ + +// pragma_struct +struct s_pragma_struct +{ +#pragma acc update /* { dg-error "expected declaration specifiers before '#pragma'" } */ +}; + +// pragma_param +void +f_pragma_param ( +#pragma acc update /* { dg-error "expected declaration specifiers before '#pragma'" } */ + void) +{ +} + +// pragma_stmt +void +f2 (void) +{ + if (0) +#pragma acc update /* { dg-error "'#pragma acc update' may only be used in compound statements before '#pragma'" } */ +} + +// pragma_compound +void +f3 (void) +{ + int i = 0; +#pragma acc update device(i) +} diff --git gcc/testsuite/c-c++-common/goacc/update-1.c gcc/testsuite/c-c++-common/goacc/update-1.c new file mode 100644 index 0000000..970fdca --- /dev/null +++ gcc/testsuite/c-c++-common/goacc/update-1.c @@ -0,0 +1,10 @@ +void +f (void) +{ +#pragma acc update /* { dg-error "'#pragma acc update' must contain at least one 'device' or 'host/self' clause" } */ + + int i = 0; +#pragma acc update device(i) +#pragma acc update host(i) +#pragma acc update self(i) +} diff --git libgomp/ChangeLog.gomp libgomp/ChangeLog.gomp index d348672..8876a2e 100644 --- libgomp/ChangeLog.gomp +++ libgomp/ChangeLog.gomp @@ -1,3 +1,9 @@ +2014-06-06 Thomas Schwinge + James Norris + + * libgomp.map (GOACC_2.0): Add GOACC_update. + * oacc-parallel.c (GOACC_update): New function. + 2014-03-18 Ilya Verbin * libgomp.map (GOMP_4.0.1): New symbol version. diff --git libgomp/libgomp.map libgomp/libgomp.map index f169f46..c575be3 100644 --- libgomp/libgomp.map +++ libgomp/libgomp.map @@ -242,4 +242,5 @@ GOACC_2.0 { GOACC_data_start; GOACC_kernels; GOACC_parallel; + GOACC_update; }; diff --git libgomp/oacc-parallel.c libgomp/oacc-parallel.c index b1fd898..79b6254 100644 --- libgomp/oacc-parallel.c +++ libgomp/oacc-parallel.c @@ -104,3 +104,27 @@ GOACC_kernels (int device, void (*fn) (void *), const void *openmp_target, GOACC_parallel (device, fn, openmp_target, mapnum, hostaddrs, sizes, kinds, num_gangs, num_workers, vector_length); } + + +void +GOACC_update (int device, const void *openmp_target, size_t mapnum, + void **hostaddrs, size_t *sizes, unsigned short *kinds) +{ + unsigned char kinds_[mapnum]; + size_t i; + + /* TODO. Eventually, we'll be interpreting all mapping kinds according to + the OpenACC semantics; for now we're re-using what is implemented for + OpenMP. */ + for (i = 0; i < mapnum; ++i) + { + unsigned char kind = kinds[i]; + unsigned char align = kinds[i] >> 8; + if (kind > 4) + gomp_fatal ("memory mapping kind %x for %zd is not yet supported", + kind, i); + + kinds_[i] = kind | align << 3; + } + GOMP_target_update (device, openmp_target, mapnum, hostaddrs, sizes, kinds_); +}