From patchwork Wed Nov 5 16:56:06 2014 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 8bit X-Patchwork-Submitter: Thomas Schwinge X-Patchwork-Id: 407090 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 194FE140098 for ; Thu, 6 Nov 2014 03:56:53 +1100 (AEDT) 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:in-reply-to:references:date:message-id :mime-version:content-type; q=dns; s=default; b=O6TOgXuphAQdXknJ BvUu810fAeMAO3JE56ucxQbYtxiyGkp4/yDFYDvduNKEOKcqgGrxgRQz/eQ45lht FRASVBukPcVtQNxpbnW7Ltug/rZVYJigd60B6vym8feCNh03tsYbVpzVjjAOLzFc KAEOzJoy7nDjSfiHmk7QKQFITGU= 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:in-reply-to:references:date:message-id :mime-version:content-type; s=default; bh=hUi3uEv8hkt6R75xHO1h0P ZA2GQ=; b=XsJnAjuApOk8HM6HfUNZgbJU1D+BEz/q9zTOioyl9mHCSg7XIVbV1C eiqX7LUfSkHrRkvewVTeZdMBaGUHvYof5eOW52uaYCHFW1Us+vY37qTTzfMJC4Wp VXhqhXey3QkqP6zKOEU/qd2ofiqPvdTx0uYNBzkk+aGEUsPDdZ7fE= Received: (qmail 8955 invoked by alias); 5 Nov 2014 16:56:37 -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 8936 invoked by uid 89); 5 Nov 2014 16:56:36 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-1.9 required=5.0 tests=BAYES_00, RCVD_IN_DNSWL_NONE, SPF_HELO_PASS autolearn=ham version=3.3.2 X-HELO: smtprelay02.ispgateway.de Received: from smtprelay02.ispgateway.de (HELO smtprelay02.ispgateway.de) (80.67.18.14) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with (AES256-GCM-SHA384 encrypted) ESMTPS; Wed, 05 Nov 2014 16:56:33 +0000 Received: from [87.149.54.146] (helo=stokes.schwinge.homeip.net) by smtprelay02.ispgateway.de with esmtpa (Exim 4.84) (envelope-from ) id 1Xm3si-0007ID-4h for gcc-patches@gcc.gnu.org; Wed, 05 Nov 2014 17:56:28 +0100 Received: (qmail 1798 invoked from network); 5 Nov 2014 16:56:09 -0000 Received: from kepler.schwinge.homeip.net (192.168.111.7) by stokes.schwinge.homeip.net with QMQP; 5 Nov 2014 16:56:09 -0000 Received: (nullmailer pid 18317 invoked by uid 1000); Wed, 05 Nov 2014 16:56:08 -0000 From: Thomas Schwinge To: "gcc-patches@gcc.gnu.org" , fortran@gcc.gnu.org Cc: Cesar Philippidis Subject: [gomp4] OpenACC update host/self maintenance (was: acc enter/exit data) In-Reply-To: <5452D398.2030709@codesourcery.com> References: <5452D398.2030709@codesourcery.com> User-Agent: Notmuch/0.9-101-g81dad07 (http://notmuchmail.org) Emacs/24.3.1 (i586-pc-linux-gnu) Date: Wed, 05 Nov 2014 17:56:06 +0100 Message-ID: <87tx2dd1sp.fsf@kepler.schwinge.homeip.net> MIME-Version: 1.0 X-Df-Sender: dGhvbWFzQHNjaHdpbmdlLm5hbWU= X-IsSubscribed: yes Hi! On Thu, 30 Oct 2014 17:11:04 -0700, Cesar Philippidis wrote: > gcc/fortran/ > * gfortran.h (enum OMP_LIST_HOST): Remove. > (enum OMP_LIST_DEVICE, OMP_LIST_DEVICE): Remove. > * dump-parse-tree.c (show_omp_clauses): Remove OMP_LIST_HOST and > OMP_LIST_DEVICE from here also. > * openmp.c (OMP_CLAUSE_SELF): New define. > (gfc_match_omp_clauses): Update handling of OMP_CLAUSE_HOST and > OMP_CLAUSE_DEVICE. Add support for OMP_CLAUSE_SELF. > * trans-openmp.c (gfc_trans_omp_clauses): Remove support for > OMP_LIST_HOST and OMP_LIST_DEVICE since they are treated as memory > maps now. > (gfc_trans_oacc_executable_directive): Remove stale EXEC_OACC_WAIT. Applied to gomp-4_0-branch in r217148: commit a7bba5ecc7c62a022616f55ff1d8fb48266fcb67 Author: tschwinge Date: Wed Nov 5 16:54:07 2014 +0000 OpenACC update host/self maintenance. gcc/c/ * c-parser.c (c_parser_omp_clause_name) <"host">: Return PRAGMA_OMP_CLAUSE_HOST. gcc/c/ (c_parser_oacc_data_clause): Group PRAGMA_OMP_CLAUSE_SELF next to PRAGMA_OMP_CLAUSE_HOST. gcc/cp/ * parser.c (cp_parser_oacc_data_clause): Group PRAGMA_OMP_CLAUSE_SELF next to PRAGMA_OMP_CLAUSE_HOST. gcc/fortran/ * openmp.c (OMP_CLAUSE_HOST, OMP_CLAUSE_SELF): Merge into the new OMP_CLAUSE_HOST_SELF. Update all users. gcc/ * tree-core.h (enum omp_clause_code): Remove OMP_CLAUSE_HOST and OMP_CLAUSE_OACC_DEVICE. Update all users. gcc/testsuite/ * c-c++-common/goacc/update-1.c: Extend. * gfortran.dg/goacc/assumed.f95: Likewise. * gfortran.dg/goacc/coarray.f95: Likewise. * gfortran.dg/goacc/cray.f95: Likewise. * gfortran.dg/goacc/literal.f95: Likewise. * gfortran.dg/goacc/parameter.f95: Likewise. libgomp/ * testsuite/libgomp.oacc-c-c++-common/update-1-2.c: New file. * testsuite/libgomp.oacc-fortran/data-4-2.f90: Likewise. * testsuite/libgomp.oacc-fortran/data-4.f90: In one instance, use the self clause instead of host clause. git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gomp-4_0-branch@217148 138bc75d-0d04-0410-961f-82ee72b054a4 --- gcc/ChangeLog.gomp | 3 + gcc/c/ChangeLog.gomp | 6 + gcc/c/c-parser.c | 10 +- gcc/cp/ChangeLog.gomp | 3 + gcc/cp/parser.c | 8 +- gcc/fortran/ChangeLog.gomp | 3 + gcc/fortran/openmp.c | 21 +- gcc/gimplify.c | 4 - gcc/omp-low.c | 4 - gcc/testsuite/ChangeLog.gomp | 7 + gcc/testsuite/c-c++-common/goacc/update-1.c | 5 + gcc/testsuite/gfortran.dg/goacc/assumed.f95 | 8 +- gcc/testsuite/gfortran.dg/goacc/coarray.f95 | 3 +- gcc/testsuite/gfortran.dg/goacc/cray.f95 | 6 +- gcc/testsuite/gfortran.dg/goacc/literal.f95 | 5 +- gcc/testsuite/gfortran.dg/goacc/parameter.f95 | 3 +- gcc/tree-core.h | 10 +- gcc/tree-pretty-print.c | 6 - gcc/tree.c | 6 - libgomp/ChangeLog.gomp | 5 + .../libgomp.oacc-c-c++-common/update-1-2.c | 282 +++++++++++++++++++++ .../{data-4.f90 => data-4-2.f90} | 8 +- libgomp/testsuite/libgomp.oacc-fortran/data-4.f90 | 2 +- 23 files changed, 353 insertions(+), 65 deletions(-) Grüße, Thomas diff --git gcc/ChangeLog.gomp gcc/ChangeLog.gomp index 2c2b349..d140a35 100644 --- gcc/ChangeLog.gomp +++ gcc/ChangeLog.gomp @@ -1,5 +1,8 @@ 2014-11-05 Thomas Schwinge + * tree-core.h (enum omp_clause_code): Remove OMP_CLAUSE_HOST and + OMP_CLAUSE_OACC_DEVICE. Update all users. + * gimplify.c (gimplify_oacc_cache): New function. (gimplify_expr): Use it for OACC_CACHE. (gimplify_scan_omp_clauses, gimplify_adjust_omp_clauses): Handle diff --git gcc/c/ChangeLog.gomp gcc/c/ChangeLog.gomp index 70278b9..a223a17 100644 --- gcc/c/ChangeLog.gomp +++ gcc/c/ChangeLog.gomp @@ -1,5 +1,11 @@ 2014-11-05 Thomas Schwinge + * c-parser.c (c_parser_omp_clause_name) <"host">: Return + PRAGMA_OMP_CLAUSE_HOST. + + * c-parser.c (c_parser_oacc_data_clause): Group + PRAGMA_OMP_CLAUSE_SELF next to PRAGMA_OMP_CLAUSE_HOST. + * c-parser.c (c_parser_oacc_cache): Generate OACC_CACHE. * c-typeck.c (c_finish_omp_clauses): Handle OMP_CLAUSE__CACHE_. diff --git gcc/c/c-parser.c gcc/c/c-parser.c index 40d4314..bd2864f 100644 --- gcc/c/c-parser.c +++ gcc/c/c-parser.c @@ -9832,7 +9832,7 @@ c_parser_omp_clause_name (c_parser *parser) break; case 'h': if (!strcmp ("host", p)) - result = PRAGMA_OMP_CLAUSE_SELF; + result = PRAGMA_OMP_CLAUSE_HOST; break; case 'i': if (!strcmp ("inbranch", p)) @@ -10187,8 +10187,6 @@ c_parser_oacc_data_clause (c_parser *parser, pragma_omp_clause c_kind, enum omp_clause_map_kind kind; switch (c_kind) { - default: - gcc_unreachable (); case PRAGMA_OMP_CLAUSE_COPY: kind = OMP_CLAUSE_MAP_FORCE_TOFROM; break; @@ -10208,6 +10206,7 @@ c_parser_oacc_data_clause (c_parser *parser, pragma_omp_clause c_kind, kind = OMP_CLAUSE_MAP_FORCE_TO; break; case PRAGMA_OMP_CLAUSE_HOST: + case PRAGMA_OMP_CLAUSE_SELF: kind = OMP_CLAUSE_MAP_FORCE_FROM; break; case PRAGMA_OMP_CLAUSE_PRESENT: @@ -10225,9 +10224,8 @@ 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; + default: + gcc_unreachable (); } tree nl, c; nl = c_parser_omp_var_list_parens (parser, OMP_CLAUSE_MAP, list); diff --git gcc/cp/ChangeLog.gomp gcc/cp/ChangeLog.gomp index 46d4912..f5d400f 100644 --- gcc/cp/ChangeLog.gomp +++ gcc/cp/ChangeLog.gomp @@ -1,5 +1,8 @@ 2014-11-05 Thomas Schwinge + * parser.c (cp_parser_oacc_data_clause): Group + PRAGMA_OMP_CLAUSE_SELF next to PRAGMA_OMP_CLAUSE_HOST. + * parser.c (cp_parser_oacc_cache): Generate OACC_CACHE. * semantics.c (finish_omp_clauses): Handle OMP_CLAUSE__CACHE_. diff --git gcc/cp/parser.c gcc/cp/parser.c index ea4ad2f..9c2c3ca 100644 --- gcc/cp/parser.c +++ gcc/cp/parser.c @@ -27812,8 +27812,6 @@ cp_parser_oacc_data_clause (cp_parser *parser, pragma_omp_clause c_kind, enum omp_clause_map_kind kind; switch (c_kind) { - default: - gcc_unreachable (); case PRAGMA_OMP_CLAUSE_COPY: kind = OMP_CLAUSE_MAP_FORCE_TOFROM; break; @@ -27833,6 +27831,7 @@ cp_parser_oacc_data_clause (cp_parser *parser, pragma_omp_clause c_kind, kind = OMP_CLAUSE_MAP_FORCE_TO; break; case PRAGMA_OMP_CLAUSE_HOST: + case PRAGMA_OMP_CLAUSE_SELF: kind = OMP_CLAUSE_MAP_FORCE_FROM; break; case PRAGMA_OMP_CLAUSE_PRESENT: @@ -27850,9 +27849,8 @@ cp_parser_oacc_data_clause (cp_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; + default: + gcc_unreachable (); } tree nl, c; nl = cp_parser_omp_var_list (parser, OMP_CLAUSE_MAP, list); diff --git gcc/fortran/ChangeLog.gomp gcc/fortran/ChangeLog.gomp index 98e3971..d10560e 100644 --- gcc/fortran/ChangeLog.gomp +++ gcc/fortran/ChangeLog.gomp @@ -1,5 +1,8 @@ 2014-11-05 Thomas Schwinge + * openmp.c (OMP_CLAUSE_HOST, OMP_CLAUSE_SELF): Merge into the new + OMP_CLAUSE_HOST_SELF. Update all users. + * gfortran.texi: Update for OpenACC. * intrinsic.texi: Likewise. * invoke.texi: Likewise. diff --git gcc/fortran/openmp.c gcc/fortran/openmp.c index c7af004..959798a 100644 --- gcc/fortran/openmp.c +++ gcc/fortran/openmp.c @@ -445,13 +445,12 @@ match_oacc_clause_gang (gfc_omp_clauses *cp) #define OMP_CLAUSE_INDEPENDENT (1ULL << 49) #define OMP_CLAUSE_USE_DEVICE (1ULL << 50) #define OMP_CLAUSE_DEVICE_RESIDENT (1ULL << 51) -#define OMP_CLAUSE_HOST (1ULL << 52) +#define OMP_CLAUSE_HOST_SELF (1ULL << 52) #define OMP_CLAUSE_OACC_DEVICE (1ULL << 53) #define OMP_CLAUSE_WAIT (1ULL << 54) #define OMP_CLAUSE_DELETE (1ULL << 55) #define OMP_CLAUSE_AUTO (1ULL << 56) #define OMP_CLAUSE_TILE (1ULL << 57) -#define OMP_CLAUSE_SELF (1ULL << 58) /* Helper function for OpenACC and OpenMP clauses involving memory mapping. */ @@ -682,24 +681,20 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, unsigned long long mask, true) == MATCH_YES) continue; - if ((mask & OMP_CLAUSE_HOST) - && gfc_match ("host ( ") == MATCH_YES - && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP], - OMP_MAP_FORCE_FROM)) - continue; if ((mask & OMP_CLAUSE_OACC_DEVICE) && gfc_match ("device ( ") == MATCH_YES && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP], OMP_MAP_FORCE_TO)) continue; + if ((mask & OMP_CLAUSE_HOST_SELF) + && (gfc_match ("host ( ") == MATCH_YES + || gfc_match ("self ( ") == MATCH_YES) + && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP], + OMP_MAP_FORCE_FROM)) + continue; if ((mask & OMP_CLAUSE_TILE) && match_oacc_expr_list ("tile (", &c->tile_list, true) == MATCH_YES) continue; - if ((mask & OMP_CLAUSE_SELF) - && gfc_match ("self ( ") == MATCH_YES - && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP], - OMP_MAP_FORCE_FROM)) - continue; if ((mask & OMP_CLAUSE_SEQ) && !c->seq && gfc_match ("seq") == MATCH_YES) { @@ -1170,7 +1165,7 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, unsigned long long mask, | OMP_CLAUSE_PRESENT_OR_COPYIN | OMP_CLAUSE_PRESENT_OR_COPYOUT \ | OMP_CLAUSE_PRESENT_OR_CREATE) #define OACC_UPDATE_CLAUSES \ - (OMP_CLAUSE_IF | OMP_CLAUSE_ASYNC | OMP_CLAUSE_HOST | OMP_CLAUSE_SELF \ + (OMP_CLAUSE_IF | OMP_CLAUSE_ASYNC | OMP_CLAUSE_HOST_SELF \ | OMP_CLAUSE_OACC_DEVICE | OMP_CLAUSE_WAIT) #define OACC_ENTER_DATA_CLAUSES \ (OMP_CLAUSE_IF | OMP_CLAUSE_ASYNC | OMP_CLAUSE_WAIT | OMP_CLAUSE_COPYIN \ diff --git gcc/gimplify.c gcc/gimplify.c index d58876f..5aeb726 100644 --- gcc/gimplify.c +++ gcc/gimplify.c @@ -6288,8 +6288,6 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, remove = true; break; - case OMP_CLAUSE_HOST: - case OMP_CLAUSE_OACC_DEVICE: case OMP_CLAUSE_DEVICE_RESIDENT: case OMP_CLAUSE_USE_DEVICE: case OMP_CLAUSE_GANG: @@ -6692,8 +6690,6 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, tree *list_p) case OMP_CLAUSE_VECTOR_LENGTH: break; - case OMP_CLAUSE_HOST: - case OMP_CLAUSE_OACC_DEVICE: case OMP_CLAUSE_DEVICE_RESIDENT: case OMP_CLAUSE_USE_DEVICE: case OMP_CLAUSE_GANG: diff --git gcc/omp-low.c gcc/omp-low.c index 1c9d942..0b45e69 100644 --- gcc/omp-low.c +++ gcc/omp-low.c @@ -1977,8 +1977,6 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) install_var_local (decl, ctx); break; - case OMP_CLAUSE_HOST: - case OMP_CLAUSE_OACC_DEVICE: case OMP_CLAUSE_DEVICE_RESIDENT: case OMP_CLAUSE_USE_DEVICE: case OMP_CLAUSE_GANG: @@ -2125,8 +2123,6 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) case OMP_CLAUSE_WAIT: break; - case OMP_CLAUSE_HOST: - case OMP_CLAUSE_OACC_DEVICE: case OMP_CLAUSE_DEVICE_RESIDENT: case OMP_CLAUSE_USE_DEVICE: case OMP_CLAUSE_GANG: diff --git gcc/testsuite/ChangeLog.gomp gcc/testsuite/ChangeLog.gomp index 1faf0fa..a02f58a 100644 --- gcc/testsuite/ChangeLog.gomp +++ gcc/testsuite/ChangeLog.gomp @@ -1,5 +1,12 @@ 2014-11-05 Thomas Schwinge + * c-c++-common/goacc/update-1.c: Extend. + * gfortran.dg/goacc/assumed.f95: Likewise. + * gfortran.dg/goacc/coarray.f95: Likewise. + * gfortran.dg/goacc/cray.f95: Likewise. + * gfortran.dg/goacc/literal.f95: Likewise. + * gfortran.dg/goacc/parameter.f95: Likewise. + * c-c++-common/goacc/cache-1.c: New file. * gfortran.dg/goacc/data-tree.f95: Remove dg-prune-output directive. diff --git gcc/testsuite/c-c++-common/goacc/update-1.c gcc/testsuite/c-c++-common/goacc/update-1.c index 2a3a910..97e9379 100644 --- gcc/testsuite/c-c++-common/goacc/update-1.c +++ gcc/testsuite/c-c++-common/goacc/update-1.c @@ -8,5 +8,10 @@ f (void) #pragma acc update device(i) #pragma acc update host(i) #pragma acc update self(i) +#pragma acc update device(a[1:3]) +#pragma acc update host(a[1:3]) +#pragma acc update self(a[1:3]) +#pragma acc update device(a(1:3)) /* { dg-error "expected '\\\)' before '\\\(' token" } */ #pragma acc update host(a(1:3)) /* { dg-error "expected '\\\)' before '\\\(' token" } */ +#pragma acc update self(a(1:3)) /* { dg-error "expected '\\\)' before '\\\(' token" } */ } diff --git gcc/testsuite/gfortran.dg/goacc/assumed.f95 gcc/testsuite/gfortran.dg/goacc/assumed.f95 index 15bfa0c..3287241 100644 --- gcc/testsuite/gfortran.dg/goacc/assumed.f95 +++ gcc/testsuite/gfortran.dg/goacc/assumed.f95 @@ -19,8 +19,9 @@ contains do i = 1,5 enddo !$acc end parallel loop - !$acc update host (a) ! { dg-error "Assumed size" } !$acc update device (a) ! { dg-error "Assumed size" } + !$acc update host (a) ! { dg-error "Assumed size" } + !$acc update self (a) ! { dg-error "Assumed size" } end subroutine assumed_size subroutine assumed_rank(a) implicit none @@ -39,7 +40,8 @@ contains do i = 1,5 enddo !$acc end parallel loop - !$acc update host (a) ! { dg-error "Assumed rank" } !$acc update device (a) ! { dg-error "Assumed rank" } + !$acc update host (a) ! { dg-error "Assumed rank" } + !$acc update self (a) ! { dg-error "Assumed rank" } end subroutine assumed_rank -end module test \ No newline at end of file +end module test diff --git gcc/testsuite/gfortran.dg/goacc/coarray.f95 gcc/testsuite/gfortran.dg/goacc/coarray.f95 index ab13157..4f1224e 100644 --- gcc/testsuite/gfortran.dg/goacc/coarray.f95 +++ gcc/testsuite/gfortran.dg/goacc/coarray.f95 @@ -27,8 +27,9 @@ contains !$acc cache (a) enddo !$acc end parallel loop - !$acc update host (a) !$acc update device (a) + !$acc update host (a) + !$acc update self (a) end subroutine oacc1 end module test ! { dg-prune-output "ACC cache unimplemented" } diff --git gcc/testsuite/gfortran.dg/goacc/cray.f95 gcc/testsuite/gfortran.dg/goacc/cray.f95 index 3225b28..8f2c077 100644 --- gcc/testsuite/gfortran.dg/goacc/cray.f95 +++ gcc/testsuite/gfortran.dg/goacc/cray.f95 @@ -28,8 +28,9 @@ contains !$acc cache (pointee) ! TODO: This must fail, as in openacc-1_0-branch enddo !$acc end parallel loop - !$acc update host (pointee) ! { dg-error "Cray pointee" } !$acc update device (pointee) ! { dg-error "Cray pointee" } + !$acc update host (pointee) ! { dg-error "Cray pointee" } + !$acc update self (pointee) ! { dg-error "Cray pointee" } !$acc data copy (ptr) !$acc end data !$acc data deviceptr (ptr) ! { dg-error "Cray pointer" } @@ -47,8 +48,9 @@ contains !$acc cache (ptr) ! TODO: This must fail, as in openacc-1_0-branch enddo !$acc end parallel loop - !$acc update host (ptr) !$acc update device (ptr) + !$acc update host (ptr) + !$acc update self (ptr) end subroutine oacc1 end module test ! { dg-prune-output "unimplemented" } diff --git gcc/testsuite/gfortran.dg/goacc/literal.f95 gcc/testsuite/gfortran.dg/goacc/literal.f95 index bdbf66d..e6760d0 100644 --- gcc/testsuite/gfortran.dg/goacc/literal.f95 +++ gcc/testsuite/gfortran.dg/goacc/literal.f95 @@ -23,7 +23,8 @@ contains !$acc cache (10) ! { dg-error "Syntax error" } enddo !$acc end parallel loop - !$acc update host (10) ! { dg-error "Syntax error" } !$acc update device (10) ! { dg-error "Syntax error" } + !$acc update host (10) ! { dg-error "Syntax error" } + !$acc update self (10) ! { dg-error "Syntax error" } end subroutine oacc1 -end module test \ No newline at end of file +end module test diff --git gcc/testsuite/gfortran.dg/goacc/parameter.f95 gcc/testsuite/gfortran.dg/goacc/parameter.f95 index 785d7f9..1364181 100644 --- gcc/testsuite/gfortran.dg/goacc/parameter.f95 +++ gcc/testsuite/gfortran.dg/goacc/parameter.f95 @@ -24,8 +24,9 @@ contains !$acc cache (a) ! TODO: This must fail, as in openacc-1_0-branch enddo !$acc end parallel loop - !$acc update host (a) ! { dg-error "not a variable" } !$acc update device (a) ! { dg-error "not a variable" } + !$acc update host (a) ! { dg-error "not a variable" } + !$acc update self (a) ! { dg-error "not a variable" } end subroutine oacc1 end module test ! { dg-prune-output "unimplemented" } diff --git gcc/tree-core.h gcc/tree-core.h index 42ad6a0..cd38861 100644 --- gcc/tree-core.h +++ gcc/tree-core.h @@ -259,8 +259,8 @@ enum omp_clause_code { OMP_CLAUSE_TO, /* OpenACC clauses: {copy, copyin, copyout, create, delete, deviceptr, - present, present_or_copy (pcopy), present_or_copyin (pcopyin), - present_or_copyout (pcopyout), present_or_create (pcreate)} + device, host (self), present, present_or_copy (pcopy), present_or_copyin + (pcopyin), present_or_copyout (pcopyout), present_or_create (pcreate)} (variable-list). OpenMP clause: map ({alloc:,to:,from:,tofrom:,}variable-list). */ @@ -270,12 +270,6 @@ enum omp_clause_code { #pragma acc cache (variable-list). */ OMP_CLAUSE__CACHE_, - /* OpenACC clause: host (variable_list). */ - OMP_CLAUSE_HOST, - - /* OpenACC clause: device (variable_list). */ - OMP_CLAUSE_OACC_DEVICE, - /* OpenACC clause: device_resident (variable_list). */ OMP_CLAUSE_DEVICE_RESIDENT, diff --git gcc/tree-pretty-print.c gcc/tree-pretty-print.c index d678f36..1458913 100644 --- gcc/tree-pretty-print.c +++ gcc/tree-pretty-print.c @@ -335,12 +335,6 @@ dump_omp_clause (pretty_printer *buffer, tree clause, int spc, int flags) case OMP_CLAUSE__LOOPTEMP_: name = "_looptemp_"; goto print_remap; - case OMP_CLAUSE_HOST: - name = "host"; - goto print_remap; - case OMP_CLAUSE_OACC_DEVICE: - name = "device"; - goto print_remap; case OMP_CLAUSE_DEVICE_RESIDENT: name = "device_resident"; goto print_remap; diff --git gcc/tree.c gcc/tree.c index f39c63f..16d156b 100644 --- gcc/tree.c +++ gcc/tree.c @@ -271,8 +271,6 @@ unsigned const char omp_clause_num_ops[] = 2, /* OMP_CLAUSE_TO */ 2, /* OMP_CLAUSE_MAP */ 2, /* OMP_CLAUSE__CACHE_ */ - 1, /* OMP_CLAUSE_HOST */ - 1, /* OMP_CLAUSE_OACC_DEVICE */ 1, /* OMP_CLAUSE_DEVICE_RESIDENT */ 1, /* OMP_CLAUSE_USE_DEVICE */ 1, /* OMP_CLAUSE_GANG */ @@ -330,8 +328,6 @@ const char * const omp_clause_code_name[] = "to", "map", "_cache_", - "host", - "device", "device_resident", "use_device", "gang", @@ -11120,8 +11116,6 @@ walk_tree_1 (tree *tp, walk_tree_fn func, void *data, case OMP_CLAUSE: switch (OMP_CLAUSE_CODE (*tp)) { - case OMP_CLAUSE_HOST: - case OMP_CLAUSE_OACC_DEVICE: case OMP_CLAUSE_DEVICE_RESIDENT: case OMP_CLAUSE_USE_DEVICE: case OMP_CLAUSE_GANG: diff --git libgomp/ChangeLog.gomp libgomp/ChangeLog.gomp index 096a2a9..d9b92fe 100644 --- libgomp/ChangeLog.gomp +++ libgomp/ChangeLog.gomp @@ -1,5 +1,10 @@ 2014-11-05 Thomas Schwinge + * testsuite/libgomp.oacc-c-c++-common/update-1-2.c: New file. + * testsuite/libgomp.oacc-fortran/data-4-2.f90: Likewise. + * testsuite/libgomp.oacc-fortran/data-4.f90: In one instance, use + the self clause instead of host clause. + * testsuite/libgomp.oacc-c/cache-1.c: Remove directives that are expected to fail, and rename the file to... * testsuite/libgomp.oacc-c-c++-common/cache-1.c: ... this. diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/update-1-2.c libgomp/testsuite/libgomp.oacc-c-c++-common/update-1-2.c new file mode 100644 index 0000000..c7e7257 --- /dev/null +++ libgomp/testsuite/libgomp.oacc-c-c++-common/update-1-2.c @@ -0,0 +1,282 @@ +/* Copy of update-1.c with self exchanged with host for #pragma acc update. */ + +/* { dg-do run } */ +/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */ + +#include +#include +#include +#include +#include + +int +main (int argc, char **argv) +{ + int N = 8; + float *a, *b, *c; + float *d_a, *d_b, *d_c; + int i; + + a = (float *) malloc (N * sizeof (float)); + b = (float *) malloc (N * sizeof (float)); + c = (float *) malloc (N * sizeof (float)); + + d_a = (float *) acc_malloc (N * sizeof (float)); + d_b = (float *) acc_malloc (N * sizeof (float)); + d_c = (float *) acc_malloc (N * sizeof (float)); + + for (i = 0; i < N; i++) + { + a[i] = 3.0; + b[i] = 0.0; + } + + acc_map_data (a, d_a, N * sizeof (float)); + acc_map_data (b, d_b, N * sizeof (float)); + acc_map_data (c, d_c, N * sizeof (float)); + +#pragma acc update device (a[0:N], b[0:N]) + +#pragma acc parallel present (a[0:N], b[0:N]) + { + int ii; + + for (ii = 0; ii < N; ii++) + b[ii] = a[ii]; + } + +#pragma acc update self (a[0:N], b[0:N]) + + for (i = 0; i < N; i++) + { + if (a[i] != 3.0) + abort (); + + if (b[i] != 3.0) + abort (); + } + + if (!acc_is_present (&a[0], (N * sizeof (float)))) + abort (); + + if (!acc_is_present (&b[0], (N * sizeof (float)))) + abort (); + + for (i = 0; i < N; i++) + { + a[i] = 5.0; + b[i] = 1.0; + } + +#pragma acc update device (a[0:N], b[0:N]) + +#pragma acc parallel present (a[0:N], b[0:N]) + { + int ii; + + for (ii = 0; ii < N; ii++) + b[ii] = a[ii]; + } + +#pragma acc update self (a[0:N], b[0:N]) + + for (i = 0; i < N; i++) + { + if (a[i] != 5.0) + abort (); + + if (b[i] != 5.0) + abort (); + } + + if (!acc_is_present (&a[0], (N * sizeof (float)))) + abort (); + + if (!acc_is_present (&b[0], (N * sizeof (float)))) + abort (); + + for (i = 0; i < N; i++) + { + a[i] = 5.0; + b[i] = 1.0; + } + +#pragma acc update device (a[0:N], b[0:N]) + +#pragma acc parallel present (a[0:N], b[0:N]) + { + int ii; + + for (ii = 0; ii < N; ii++) + b[ii] = a[ii]; + } + +#pragma acc update host (a[0:N], b[0:N]) + + for (i = 0; i < N; i++) + { + if (a[i] != 5.0) + abort (); + + if (b[i] != 5.0) + abort (); + } + + if (!acc_is_present (&a[0], (N * sizeof (float)))) + abort (); + + if (!acc_is_present (&b[0], (N * sizeof (float)))) + abort (); + + for (i = 0; i < N; i++) + { + a[i] = 6.0; + b[i] = 0.0; + } + +#pragma acc update device (a[0:N], b[0:N]) + + for (i = 0; i < N; i++) + { + a[i] = 9.0; + } + +#pragma acc parallel present (a[0:N], b[0:N]) + { + int ii; + + for (ii = 0; ii < N; ii++) + b[ii] = a[ii]; + } + +#pragma acc update self (a[0:N], b[0:N]) + + for (i = 0; i < N; i++) + { + if (a[i] != 6.0) + abort (); + + if (b[i] != 6.0) + abort (); + } + + if (!acc_is_present (&a[0], (N * sizeof (float)))) + abort (); + + if (!acc_is_present (&b[0], (N * sizeof (float)))) + abort (); + + for (i = 0; i < N; i++) + { + a[i] = 7.0; + b[i] = 2.0; + } + +#pragma acc update device (a[0:N], b[0:N]) + + for (i = 0; i < N; i++) + { + a[i] = 9.0; + } + +#pragma acc parallel present (a[0:N], b[0:N]) + { + int ii; + + for (ii = 0; ii < N; ii++) + b[ii] = a[ii]; + } + +#pragma acc update self (a[0:N], b[0:N]) + + for (i = 0; i < N; i++) + { + if (a[i] != 7.0) + abort (); + + if (b[i] != 7.0) + abort (); + } + + for (i = 0; i < N; i++) + { + a[i] = 9.0; + } + +#pragma acc update device (a[0:N]) + +#pragma acc parallel present (a[0:N], b[0:N]) + { + int ii; + + for (ii = 0; ii < N; ii++) + b[ii] = a[ii]; + } + +#pragma acc update self (a[0:N], b[0:N]) + + for (i = 0; i < N; i++) + { + if (a[i] != 9.0) + abort (); + + if (b[i] != 9.0) + abort (); + } + + if (!acc_is_present (&a[0], (N * sizeof (float)))) + abort (); + + if (!acc_is_present (&b[0], (N * sizeof (float)))) + abort (); + + for (i = 0; i < N; i++) + { + a[i] = 5.0; + } + +#pragma acc update device (a[0:N]) + + for (i = 0; i < N; i++) + { + a[i] = 6.0; + } + +#pragma acc update device (a[0:N >> 1]) + +#pragma acc parallel present (a[0:N], b[0:N]) + { + int ii; + + for (ii = 0; ii < N; ii++) + b[ii] = a[ii]; + } + +#pragma acc update self (a[0:N], b[0:N]) + + for (i = 0; i < (N >> 1); i++) + { + if (a[i] != 6.0) + abort (); + + if (b[i] != 6.0) + abort (); + } + + for (i = (N >> 1); i < N; i++) + { + if (a[i] != 5.0) + abort (); + + if (b[i] != 5.0) + abort (); + } + + if (!acc_is_present (&a[0], (N * sizeof (float)))) + abort (); + + if (!acc_is_present (&b[0], (N * sizeof (float)))) + abort (); + + return 0; +} diff --git libgomp/testsuite/libgomp.oacc-fortran/data-4.f90 libgomp/testsuite/libgomp.oacc-fortran/data-4-2.f90 similarity index 91% copy from libgomp/testsuite/libgomp.oacc-fortran/data-4.f90 copy to libgomp/testsuite/libgomp.oacc-fortran/data-4-2.f90 index 41c45fb..16a8598 100644 --- libgomp/testsuite/libgomp.oacc-fortran/data-4.f90 +++ libgomp/testsuite/libgomp.oacc-fortran/data-4-2.f90 @@ -1,3 +1,5 @@ +! Copy of data-4.f90 with self exchanged with host for !acc update. + ! { dg-do run } program asyncwait @@ -24,7 +26,7 @@ program asyncwait end do !$acc end parallel - !$acc update host (a(1:N), b(1:N)) async wait + !$acc update self (a(1:N), b(1:N)) async wait !$acc wait do i = 1, N @@ -78,7 +80,7 @@ program asyncwait end do !$acc end parallel - !$acc update host (a(1:N), b(1:N), c(1:N), d(1:N)) async (1) wait (1) + !$acc update self (a(1:N), b(1:N), c(1:N), d(1:N)) async (1) wait (1) !$acc wait (1) @@ -122,7 +124,7 @@ program asyncwait end do !$acc end parallel - !$acc update host (a(1:N), b(1:N), c(1:N), d(1:N), e(1:N)) async (1) wait (1) + !$acc update self (a(1:N), b(1:N), c(1:N), d(1:N), e(1:N)) async (1) wait (1) !$acc wait (1) !$acc exit data delete (N, a(1:N), b(1:N), c(1:N), d(1:N), e(1:N)) diff --git libgomp/testsuite/libgomp.oacc-fortran/data-4.f90 libgomp/testsuite/libgomp.oacc-fortran/data-4.f90 index 41c45fb..f6886b0 100644 --- libgomp/testsuite/libgomp.oacc-fortran/data-4.f90 +++ libgomp/testsuite/libgomp.oacc-fortran/data-4.f90 @@ -44,7 +44,7 @@ program asyncwait end do !$acc end parallel - !$acc update host (a(1:N), b(1:N)) async (1) wait (1) + !$acc update self (a(1:N), b(1:N)) async (1) wait (1) !$acc wait (1) do i = 1, N