From patchwork Thu Jun 5 14:00:16 2014 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Thomas Schwinge X-Patchwork-Id: 356421 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 CD8EC1400E9 for ; Fri, 6 Jun 2014 00:00:46 +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:subject:date:message-id:in-reply-to:references:mime-version :content-type; q=dns; s=default; b=YGBRIs7UHhuxDx72oUMNKN2bV57E+ DQ9L/K8TVOxlz3J6TSVqlEdk3YgTYUX6gMV1tUCMwUjIEkoKLSL4xJmOWlsPg+BG h5QBMsSf79AMAMKXxcvLhBob5rGa+gqNn3VJ2xzyPAh3Zkxg5546gdyUNs2S+82g AiuToJcysvu95k= 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:subject:date:message-id:in-reply-to:references:mime-version :content-type; s=default; bh=EiOAi7pSpi2qa+DOFVlruCUDg3Q=; b=tQO L5K/zd/tuI1F2+OunQi0PjIhVg6NBDvCW1p7dDxenyehxGupnky+p3VsDUTJbPEB dB8gB80IidKdFVFUr8v0/6VThPlfU1jNaTFSBXkiVpkqgWxzmi8g3uiNE7/SgWWx 9D9Mxnq+FVBpBEwPA/VJfUC0fslo/hYCTp1NIBb8= Received: (qmail 32202 invoked by alias); 5 Jun 2014 14:00:31 -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 32136 invoked by uid 89); 5 Jun 2014 14:00:30 -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; Thu, 05 Jun 2014 14:00:25 +0000 Received: from svr-orw-exc-10.mgc.mentorg.com ([147.34.98.58]) by relay1.mentorg.com with esmtp id 1WsYDN-00073G-Nj from Thomas_Schwinge@mentor.com for gcc-patches@gcc.gnu.org; Thu, 05 Jun 2014 07:00:21 -0700 Received: from SVR-IES-FEM-02.mgc.mentorg.com ([137.202.0.106]) by SVR-ORW-EXC-10.mgc.mentorg.com with Microsoft SMTPSVC(6.0.3790.4675); Thu, 5 Jun 2014 07:00:21 -0700 Received: from feldtkeller.schwinge.homeip.net (137.202.0.76) by SVR-IES-FEM-02.mgc.mentorg.com (137.202.0.106) with Microsoft SMTP Server id 14.2.247.3; Thu, 5 Jun 2014 15:00:20 +0100 From: Thomas Schwinge To: Subject: [GOMP4, COMMITTED] OpenACC deviceptr clause. Date: Thu, 5 Jun 2014 16:00:16 +0200 Message-ID: <1401976816-10577-1-git-send-email-thomas@codesourcery.com> In-Reply-To: <87ppnuvbv6.fsf@schwinge.name> References: <87ppnuvbv6.fsf@schwinge.name> MIME-Version: 1.0 From: tschwinge gcc/c/ * c-typeck.c (handle_omp_array_sections, c_finish_omp_clauses): Handle OMP_CLAUSE_MAP_FORCE_DEVICEPTR. gcc/ * gimplify.c (gimplify_scan_omp_clauses) (gimplify_adjust_omp_clauses): Handle OMP_CLAUSE_MAP_FORCE_DEVICEPTR. * omp-low.c (scan_sharing_clauses, lower_oacc_offload) (lower_omp_target): Likewise. * tree-core.h (enum omp_clause_map_kind) : Update comment. gcc/testsuite/ * c-c++-common/goacc/data-clause-duplicate-1.c: The OpenACC deviceptr clause is now supported. * c-c++-common/goacc/deviceptr-1.c: Extend. * c-c++-common/goacc/deviceptr-2.c: New file. git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gomp-4_0-branch@211278 138bc75d-0d04-0410-961f-82ee72b054a4 --- gcc/ChangeLog.gomp | 8 +++ gcc/c/ChangeLog.gomp | 5 ++ gcc/c/c-typeck.c | 5 +- gcc/gimplify.c | 7 ++- gcc/omp-low.c | 60 +++++++++++++++++++--- gcc/testsuite/ChangeLog.gomp | 5 ++ .../c-c++-common/goacc/data-clause-duplicate-1.c | 4 +- gcc/testsuite/c-c++-common/goacc/deviceptr-1.c | 22 +++++++- gcc/testsuite/c-c++-common/goacc/deviceptr-2.c | 23 +++++++++ gcc/tree-core.h | 3 +- 10 files changed, 127 insertions(+), 15 deletions(-) create mode 100644 gcc/testsuite/c-c++-common/goacc/deviceptr-2.c diff --git gcc/ChangeLog.gomp gcc/ChangeLog.gomp index 7371aa5..88f09b3 100644 --- gcc/ChangeLog.gomp +++ gcc/ChangeLog.gomp @@ -1,5 +1,13 @@ 2014-06-05 Thomas Schwinge + * gimplify.c (gimplify_scan_omp_clauses) + (gimplify_adjust_omp_clauses): Handle + OMP_CLAUSE_MAP_FORCE_DEVICEPTR. + * omp-low.c (scan_sharing_clauses, lower_oacc_offload) + (lower_omp_target): Likewise. + * tree-core.h (enum omp_clause_map_kind) + : Update comment. + * gimplify.c (gimplify_scan_omp_clauses) : Don't block OMP_CLAUSE_MAP_FORCE_PRESENT. diff --git gcc/c/ChangeLog.gomp gcc/c/ChangeLog.gomp index 91978db..1e80031 100644 --- gcc/c/ChangeLog.gomp +++ gcc/c/ChangeLog.gomp @@ -1,3 +1,8 @@ +2014-06-05 Thomas Schwinge + + * c-typeck.c (handle_omp_array_sections, c_finish_omp_clauses): + Handle OMP_CLAUSE_MAP_FORCE_DEVICEPTR. + 2014-03-20 Thomas Schwinge * c-parser.c: Update comments. diff --git gcc/c/c-typeck.c gcc/c/c-typeck.c index c4ba531..839cdf7 100644 --- gcc/c/c-typeck.c +++ gcc/c/c-typeck.c @@ -11747,6 +11747,7 @@ handle_omp_array_sections (tree c) OMP_CLAUSE_SIZE (c) = size; if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP) return false; + gcc_assert (OMP_CLAUSE_MAP_KIND (c) != OMP_CLAUSE_MAP_FORCE_DEVICEPTR); tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP); OMP_CLAUSE_MAP_KIND (c2) = OMP_CLAUSE_MAP_POINTER; if (!c_mark_addressable (t)) @@ -12168,7 +12169,9 @@ c_finish_omp_clauses (tree clauses) else if (!c_mark_addressable (t)) remove = true; else if (!(OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP - && OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_POINTER) + && (OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_POINTER + || (OMP_CLAUSE_MAP_KIND (c) + == OMP_CLAUSE_MAP_FORCE_DEVICEPTR))) && !lang_hooks.types.omp_mappable_type (TREE_TYPE (t))) { error_at (OMP_CLAUSE_LOCATION (c), diff --git gcc/gimplify.c gcc/gimplify.c index 6eaf6fd..a1b6be6 100644 --- gcc/gimplify.c +++ gcc/gimplify.c @@ -6015,7 +6015,6 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, switch (OMP_CLAUSE_MAP_KIND (c)) { case OMP_CLAUSE_MAP_FORCE_DEALLOC: - case OMP_CLAUSE_MAP_FORCE_DEVICEPTR: input_location = OMP_CLAUSE_LOCATION (c); /* TODO. */ sorry ("data clause not yet implemented"); @@ -6533,6 +6532,12 @@ gimplify_adjust_omp_clauses (tree *list_p) && TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST && OMP_CLAUSE_MAP_KIND (c) != OMP_CLAUSE_MAP_POINTER) { + /* For OMP_CLAUSE_MAP_FORCE_DEVICEPTR, we'll never enter here, + because for these, TREE_CODE (DECL_SIZE (decl)) will always be + INTEGER_CST. */ + gcc_assert (OMP_CLAUSE_MAP_KIND (c) + != OMP_CLAUSE_MAP_FORCE_DEVICEPTR); + tree decl2 = DECL_VALUE_EXPR (decl); gcc_assert (TREE_CODE (decl2) == INDIRECT_REF); decl2 = TREE_OPERAND (decl2, 0); diff --git gcc/omp-low.c gcc/omp-low.c index 3e282c0..39f0598 100644 --- gcc/omp-low.c +++ gcc/omp-low.c @@ -1708,6 +1708,18 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) && !POINTER_TYPE_P (TREE_TYPE (decl))) break; } +#if 0 + /* In target regions that are not offloaded, libgomp won't pay + attention to OMP_CLAUSE_MAP_FORCE_DEVICEPTR -- but I think we need + to handle it here anyway, in order to create a visible copy of the + variable. */ + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP + && OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_FORCE_DEVICEPTR) + { + if (!is_gimple_omp_offloaded (ctx->stmt)) + break; + } +#endif if (DECL_P (decl)) { if (DECL_SIZE (decl) @@ -1723,6 +1735,10 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) } else { + gcc_assert (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP + || (OMP_CLAUSE_MAP_KIND (c) + != OMP_CLAUSE_MAP_FORCE_DEVICEPTR) + || TREE_CODE (TREE_TYPE (decl)) != ARRAY_TYPE); if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP && OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_POINTER && !OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c) @@ -1738,6 +1754,10 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) { tree base = get_base_address (decl); tree nc = OMP_CLAUSE_CHAIN (c); + gcc_assert (nc == NULL_TREE + || OMP_CLAUSE_CODE (nc) != OMP_CLAUSE_MAP + || (OMP_CLAUSE_MAP_KIND (nc) + != OMP_CLAUSE_MAP_FORCE_DEVICEPTR)); if (DECL_P (base) && nc != NULL_TREE && OMP_CLAUSE_CODE (nc) == OMP_CLAUSE_MAP @@ -1867,6 +1887,9 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) } if (DECL_P (decl)) { + gcc_assert ((OMP_CLAUSE_MAP_KIND (c) + != OMP_CLAUSE_MAP_FORCE_DEVICEPTR) + || TREE_CODE (TREE_TYPE (decl)) != ARRAY_TYPE); if (OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_POINTER && TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE && !COMPLETE_TYPE_P (TREE_TYPE (decl))) @@ -1878,6 +1901,9 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) else if (DECL_SIZE (decl) && TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST) { + gcc_assert (OMP_CLAUSE_MAP_KIND (c) + != OMP_CLAUSE_MAP_FORCE_DEVICEPTR); + tree decl2 = DECL_VALUE_EXPR (decl); gcc_assert (TREE_CODE (decl2) == INDIRECT_REF); decl2 = TREE_OPERAND (decl2, 0); @@ -9100,6 +9126,10 @@ lower_oacc_offload (gimple_stmt_iterator *gsi_p, omp_context *ctx) { x = build_receiver_ref (var, true, ctx); tree new_var = lookup_decl (var, ctx); + gcc_assert (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP + || (OMP_CLAUSE_MAP_KIND (c) + != OMP_CLAUSE_MAP_FORCE_DEVICEPTR) + || TREE_CODE (TREE_TYPE (var)) != ARRAY_TYPE); if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP && OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_POINTER && !OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c) @@ -9199,6 +9229,10 @@ lower_oacc_offload (gimple_stmt_iterator *gsi_p, omp_context *ctx) { tree var = lookup_decl_in_outer_ctx (ovar, ctx); tree x = build_sender_ref (ovar, ctx); + gcc_assert (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP + || (OMP_CLAUSE_MAP_KIND (c) + != OMP_CLAUSE_MAP_FORCE_DEVICEPTR) + || TREE_CODE (TREE_TYPE (ovar)) != ARRAY_TYPE); if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP && OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_POINTER && !OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c) @@ -9219,12 +9253,14 @@ lower_oacc_offload (gimple_stmt_iterator *gsi_p, omp_context *ctx) = OMP_CLAUSE_MAP_KIND (c); if ((!(map_kind & OMP_CLAUSE_MAP_SPECIAL) && (map_kind & OMP_CLAUSE_MAP_TO)) - || map_kind == OMP_CLAUSE_MAP_POINTER) + || map_kind == OMP_CLAUSE_MAP_POINTER + || map_kind == OMP_CLAUSE_MAP_FORCE_DEVICEPTR) gimplify_assign (avar, var, &ilist); avar = build_fold_addr_expr (avar); gimplify_assign (x, avar, &ilist); - if ((!(map_kind & OMP_CLAUSE_MAP_SPECIAL) - && (map_kind & OMP_CLAUSE_MAP_FROM)) + if (((!(map_kind & OMP_CLAUSE_MAP_SPECIAL) + && (map_kind & OMP_CLAUSE_MAP_FROM)) + || map_kind == OMP_CLAUSE_MAP_FORCE_DEVICEPTR) && !TYPE_READONLY (TREE_TYPE (var))) { x = build_sender_ref (ovar, ctx); @@ -10606,6 +10642,10 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) { x = build_receiver_ref (var, true, ctx); tree new_var = lookup_decl (var, ctx); + gcc_assert (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP + || (OMP_CLAUSE_MAP_KIND (c) + != OMP_CLAUSE_MAP_FORCE_DEVICEPTR) + || TREE_CODE (TREE_TYPE (var)) != ARRAY_TYPE); if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP && OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_POINTER && !OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c) @@ -10732,12 +10772,15 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) { tree var = lookup_decl_in_outer_ctx (ovar, ctx); tree x = build_sender_ref (ovar, ctx); + gcc_assert (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP + || (OMP_CLAUSE_MAP_KIND (c) + != OMP_CLAUSE_MAP_FORCE_DEVICEPTR) + || TREE_CODE (TREE_TYPE (ovar)) != ARRAY_TYPE); if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP && OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_POINTER && !OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c) && TREE_CODE (TREE_TYPE (ovar)) == ARRAY_TYPE) { - gcc_assert (kind == GF_OMP_TARGET_KIND_REGION); tree avar = create_tmp_var (TREE_TYPE (TREE_TYPE (x)), NULL); mark_addressable (avar); @@ -10747,19 +10790,20 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) } else if (is_gimple_reg (var)) { - gcc_assert (kind == GF_OMP_TARGET_KIND_REGION); tree avar = create_tmp_var (TREE_TYPE (var), NULL); mark_addressable (avar); enum omp_clause_map_kind map_kind = OMP_CLAUSE_MAP_KIND (c); if ((!(map_kind & OMP_CLAUSE_MAP_SPECIAL) && (map_kind & OMP_CLAUSE_MAP_TO)) - || map_kind == OMP_CLAUSE_MAP_POINTER) + || map_kind == OMP_CLAUSE_MAP_POINTER + || map_kind == OMP_CLAUSE_MAP_FORCE_DEVICEPTR) gimplify_assign (avar, var, &ilist); avar = build_fold_addr_expr (avar); gimplify_assign (x, avar, &ilist); - if ((!(map_kind & OMP_CLAUSE_MAP_SPECIAL) - && (map_kind & OMP_CLAUSE_MAP_FROM)) + if (((!(map_kind & OMP_CLAUSE_MAP_SPECIAL) + && (map_kind & OMP_CLAUSE_MAP_FROM)) + || map_kind == OMP_CLAUSE_MAP_FORCE_DEVICEPTR) && !TYPE_READONLY (TREE_TYPE (var))) { x = build_sender_ref (ovar, ctx); diff --git gcc/testsuite/ChangeLog.gomp gcc/testsuite/ChangeLog.gomp index 4e0ee28..08ec907 100644 --- gcc/testsuite/ChangeLog.gomp +++ gcc/testsuite/ChangeLog.gomp @@ -1,5 +1,10 @@ 2014-06-05 Thomas Schwinge + * c-c++-common/goacc/data-clause-duplicate-1.c: The OpenACC + deviceptr clause is now supported. + * c-c++-common/goacc/deviceptr-1.c: Extend. + * c-c++-common/goacc/deviceptr-2.c: New file. + * c-c++-common/goacc/data-clause-duplicate-1.c: Extend. * c-c++-common/goacc/present-1.c: New file. diff --git gcc/testsuite/c-c++-common/goacc/data-clause-duplicate-1.c gcc/testsuite/c-c++-common/goacc/data-clause-duplicate-1.c index 5c5ab02..7a1cf68 100644 --- gcc/testsuite/c-c++-common/goacc/data-clause-duplicate-1.c +++ gcc/testsuite/c-c++-common/goacc/data-clause-duplicate-1.c @@ -6,9 +6,7 @@ fun (void) ; #pragma acc kernels present_or_copyin(fp[3]) present_or_copyout(fp[7:4]) /* { dg-error "'fp' appears more than once in map clauses" } */ ; -#pragma acc data create(fp[:10]) deviceptr(fp) - /* { dg-error "'fp' appears more than once in map clauses" "" { target *-*-* } 9 } */ - /* { dg-message "sorry, unimplemented: data clause not yet implemented" "" { target *-*-* } 9 } */ +#pragma acc data create(fp[:10]) deviceptr(fp) /* { dg-error "'fp' appears more than once in map clauses" } */ ; #pragma acc data create(fp) present(fp) /* { dg-error "'fp' appears more than once in map clauses" } */ ; diff --git gcc/testsuite/c-c++-common/goacc/deviceptr-1.c gcc/testsuite/c-c++-common/goacc/deviceptr-1.c index 1ac63bd..cf2d809 100644 --- gcc/testsuite/c-c++-common/goacc/deviceptr-1.c +++ gcc/testsuite/c-c++-common/goacc/deviceptr-1.c @@ -61,4 +61,24 @@ fun3 (void) ; } -/* { dg-prune-output "sorry, unimplemented: data clause not yet implemented" } */ +extern struct s s1; +extern struct s s2[1]; /* { dg-error "array type has incomplete element type" "" { target c } } */ + +void +fun4 (void) +{ + struct s *s1_p = &s1; + struct s *s2_p = &s2; + +#pragma acc parallel deviceptr(s1) /* { dg-error "'s1' is not a pointer variable" } */ + ; + +#pragma acc parallel deviceptr(s2) + ; + +#pragma acc parallel deviceptr(s1_p) + s1_p = 0; + +#pragma acc parallel deviceptr(s2_p) + s2_p = 0; +} diff --git gcc/testsuite/c-c++-common/goacc/deviceptr-2.c gcc/testsuite/c-c++-common/goacc/deviceptr-2.c new file mode 100644 index 0000000..ac162b4 --- /dev/null +++ gcc/testsuite/c-c++-common/goacc/deviceptr-2.c @@ -0,0 +1,23 @@ +void +fun1 (void) +{ + char *a = 0; + +#pragma acc data deviceptr(a) + ++a; + +#pragma acc data deviceptr(a) +#pragma acc parallel + ++a; + +#pragma acc data deviceptr(a) +#pragma acc parallel deviceptr(a) + ++a; + +#pragma acc data +#pragma acc parallel deviceptr(a) + ++a; + +#pragma acc parallel deviceptr(a) + ++a; +} diff --git gcc/tree-core.h gcc/tree-core.h index 8603553..8b70c5b 100644 --- gcc/tree-core.h +++ gcc/tree-core.h @@ -1225,7 +1225,8 @@ enum omp_clause_map_kind OMP_CLAUSE_MAP_FORCE_PRESENT = OMP_CLAUSE_MAP_FORCE | OMP_CLAUSE_MAP_SPECIAL, /* Deallocate a mapping, without copying from device. */ OMP_CLAUSE_MAP_FORCE_DEALLOC, - /* Is a device pointer. */ + /* Is a device pointer. OMP_CLAUSE_SIZE for these is unused; is implicitly + POINTER_SIZE / BITS_PER_UNIT. */ OMP_CLAUSE_MAP_FORCE_DEVICEPTR, /* End marker. */