From patchwork Sun Oct 6 22:32:37 2019 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Julian Brown X-Patchwork-Id: 1172587 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@bilbo.ozlabs.org Authentication-Results: ozlabs.org; spf=pass (mailfrom) smtp.mailfrom=gcc.gnu.org (client-ip=209.132.180.131; helo=sourceware.org; envelope-from=gcc-patches-return-510373-incoming=patchwork.ozlabs.org@gcc.gnu.org; receiver=) Authentication-Results: ozlabs.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: ozlabs.org; dkim=pass (1024-bit key; unprotected) header.d=gcc.gnu.org header.i=@gcc.gnu.org header.b="ntRKYo0d"; dkim-atps=neutral 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 46mdh110jgz9sP7 for ; Mon, 7 Oct 2019 09:34:08 +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:date:message-id:in-reply-to:references :mime-version:content-transfer-encoding:content-type; q=dns; s= default; b=IrZvfaRFARuTnFdBNEow7QSb5T1u5zQ74F+TJ7bM8ZAPj/5dpzqzr wAuV3ptXGJYBw4WF8n0giWzyxHBlXC+p6/6JiVUGlAqScfauq6AKD21freLTimHz lrFVFf7XL2UhU5pyFv/8NTaJmEf5KKRXtcN578e1Awpc8r/GHy+PAk= 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:in-reply-to:references :mime-version:content-transfer-encoding:content-type; s=default; bh=rMVspt7Ee/WJe170BgKSa+1U4Ys=; b=ntRKYo0dLFhEnBcaI526J6CwxYtk +gUauU4u22C4Mjl48LL1JWiRiUBBW+Ii3ZhS3f+uI4ipOy1QALPW1D+sDA8Feh6J AFBSQUhyQa08sQ3LtaHIg130cE2nci9moc5FVvt5DTHaYWeTnHdJ5EoNFlckClDV 8w9rPQXJX/GIgjk= Received: (qmail 24695 invoked by alias); 6 Oct 2019 22:33: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 24473 invoked by uid 89); 6 Oct 2019 22:33:35 -0000 Authentication-Results: sourceware.org; auth=none X-Spam-SWARE-Status: No, score=-22.9 required=5.0 tests=AWL, BAYES_00, GIT_PATCH_0, GIT_PATCH_1, GIT_PATCH_2, GIT_PATCH_3, KAM_SHORT, SPF_PASS autolearn=ham version=3.3.1 spammy=belonging, covered, 5n, thr X-HELO: esa3.mentor.iphmx.com Received: from esa3.mentor.iphmx.com (HELO esa3.mentor.iphmx.com) (68.232.137.180) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with ESMTP; Sun, 06 Oct 2019 22:33:08 +0000 IronPort-SDR: /RXwHl+Kl+S75pYU8EIfXcXzQUPIlGr3+3s3b4LsSXo8xUhgN3v8gO/3CxdwQ7rf/p5OLQ2gMI 9CWqp8uRzeytRjM/xWStLVxuYEyhKnwIt4nbKAPc6slDostBM/1l8Rl91kQOeFBiELkKW0PvT0 KE2phfmy/L1+p8q1f0oqI3I2jm9jUbdiCbfSzeM5EJwfBjChCvYb5gZ05Mf+x5pCEDmStd22yU pgI6dm3MTydjp0zcyAiXhSsXwYLPxQVsaYVx2JIyFWGjQIx6T31m2YKPq7fm1HySBKvchywMzZ IiA= Received: from orw-gwy-01-in.mentorg.com ([192.94.38.165]) by esa3.mentor.iphmx.com with ESMTP; 06 Oct 2019 14:33:07 -0800 IronPort-SDR: /M9TwEQNvyPcjhPny8I/kxVo7HwUG8WDP8uOdRr8Z8YC7rbeYQTOz8aaZAS0rECR9W4Xew9bUh gVGwYXmKtaBD0RaEovZRMamEWB5asHqsUA68QcpvOkok0LgyMAc9vLGZiXvAJbLMH8oLLkevzV UksCtgWHx9azZN2G1kNBOHkuLZKo8mA34AeWIHg6nhNLeJtOZBPAhVc+pLah+2FbasS5NwW2JK AVj1tSYVsxHyQsFvx8YFwh3Ll6gA9ZAJt4UWagspNV/ACPa2ivDBXTCgLXjSDhLQcmO3mubEPC zTQ= From: Julian Brown To: CC: , Subject: [PATCH 4/4] OpenACC 2.6 manual deep copy support (attach/detach) Date: Sun, 6 Oct 2019 15:32:37 -0700 Message-ID: <20191006223237.81842-5-julian@codesourcery.com> In-Reply-To: <20191006223237.81842-1-julian@codesourcery.com> References: <20191006223237.81842-1-julian@codesourcery.com> MIME-Version: 1.0 X-IsSubscribed: yes This patch provides support for OpenACC 2.6 manual deep copy, without the reference-counting changes included in the previous version of the patch: https://gcc.gnu.org/ml/gcc-patches/2018-12/msg01084.html Relative to that version of the patch, I've introduced a GOMP_MAP_ATTACH_DETACH mapping type (only used internally in GCC, not exposed to libgomp) to distinguish the OpenACC semantics from the previously-used GOMP_MAP_ALWAYS_POINTER mapping, and reimplementing the Fortran derived-type component mapping to fix some bugs, improve code re-use, prepare for "class" support (a forthcoming patch), and to keep better separation between OpenACC-specific code and OpenMP. This version of the patch also includes several new tests. OK for trunk? Thanks, Julian ChangeLog gcc/c-family/ * c-pragma.h (pragma_omp_clause): Add PRAGMA_OACC_CLAUSE_ATTACH and PRAGMA_OACC_CLAUSE_DETACH. gcc/c/ * c-parser.c (c_parser_omp_clause_name): Add parsing of attach and detach clauses. (c_parser_omp_variable_list): Add ALLOW_DEREF optional parameter. Allow deref (->) in variable lists if true. (c_parser_omp_var_list_parens): Add ALLOW_DEREF optional parameter. Pass to c_parser_omp_variable_list. (c_parser_oacc_data_clause): Support attach and detach clauses. Update call to c_parser_omp_variable_list. (c_parser_oacc_all_clauses): Support attach and detach clauses. (OACC_DATA_CLAUSE_MASK, OACC_ENTER_DATA_CLAUSE_MASK) (OACC_KERNELS_CLAUSE_MASK, OACC_PARALLEL_CLAUSE_MASK): Add PRAGMA_OACC_CLAUSE_ATTACH. (OACC_EXIT_DATA_CLAUSE_MASK): Add PRAGMA_OACC_CLAUSE_DETACH. * c-typeck.c (handle_omp_array_sections_1): Reject subarrays for attach and detach. Support deref. (handle_omp_array_sections): Use GOMP_MAP_ATTACH_DETACH instead of GOMP_MAP_ALWAYS_POINTER for OpenACC. (c_oacc_check_attachments): New function. (c_finish_omp_clauses): Check attach/detach arguments for being pointers using above. Support deref. gcc/cp/ * parser.c (cp_parser_omp_clause_name): Support attach and detach clauses. (cp_parser_omp_var_list_no_open): Add ALLOW_DEREF optional parameter. Parse deref if true. (cp_parser_omp_var_list): Add ALLOW_DEREF optional parameter. Pass to cp_parser_omp_var_list_no_open. (cp_parser_oacc_data_clause): Support attach and detach clauses. Update call to cp_parser_omp_var_list_no_open. (cp_parser_oacc_all_clauses): Support attach and detach. (OACC_DATA_CLAUSE_MASK, OACC_ENTER_DATA_CLAUSE_MASK) (OACC_KERNELS_CLAUSE_MASK, OACC_PARALLEL_CLAUSE_MASK): Add PRAGMA_OACC_CLAUSE_ATTACH. (OACC_EXIT_DATA_CLAUSE_MASK): Add PRAGMA_OACC_CLAUSE_DETACH. * semantics.c (handle_omp_array_sections_1): Reject subarrays for attach and detach. (handle_omp_array_sections): Use GOMP_MAP_ATTACH_DETACH instead of GOMP_MAP_ALWAYS_POINTER for OpenACC. (cp_oacc_check_attachments): New function. (finish_omp_clauses): Use above function. Allow structure fields and class members to appear in OpenACC data clauses. Support GOMP_MAP_ATTACH_DETACH. Support deref. gcc/fortran/ * gfortran.h (gfc_omp_map_op): Add OMP_MAP_ATTACH, OMP_MAP_DETACH. * openmp.c (gfc_match_omp_variable_list): Add allow_derived parameter. Parse derived-type member accesses if true. (omp_mask2): Add OMP_CLAUSE_ATTACH and OMP_CLAUSE_DETACH. (gfc_match_omp_map_clause): Add allow_derived parameter. Pass to gfc_match_omp_variable_list. (gfc_match_omp_clauses): Support attach and detach. Support derived types for appropriate OpenACC directives. (OACC_PARALLEL_CLAUSES, OACC_KERNELS_CLAUSES, OACC_DATA_CLAUSES) (OACC_ENTER_DATA_CLAUSES): Add OMP_CLAUSE_ATTACH. (OACC_EXIT_DATA_CLAUSES): Add OMP_CLAUSE_DETACH. (check_symbol_not_pointer): Don't disallow pointer objects of derived type. (resolve_oacc_data_clauses): Don't disallow allocatable derived types. (resolve_omp_clauses): Perform duplicate checking only for non-derived type component accesses (plain variables and arrays or array sections). Support component refs. * trans-expr.c (gfc_conv_component_ref, conv_parent_component_references): Make global. (gfc_auto_dereference_var): New function, broken out of... (gfc_conv_variable): ...here. Call above function. * trans-openmp.c (gfc_omp_privatize_by_reference): Support component refs. (gfc_trans_omp_array_section): New function, broken out of... (gfc_trans_omp_clauses): ...here. Support component refs/derived types, attach and detach clauses. * trans.h (gfc_conv_component_ref, conv_parent_component_references, gfc_auto_dereference_var): Add prototypes. gcc/ * gimplify.c (gimplify_omp_var_data): Add GOVD_MAP_HAS_ATTACHMENTS. (insert_struct_comp_map): Support derived-type member mappings for arrays with descriptors which use GOMP_MAP_TO_PSET. Support GOMP_MAP_ATTACH_DETACH. (gimplify_scan_omp_clauses): Tidy up OACC_ENTER_DATA/OACC_EXIT_DATA mappings. Handle attach/detach clauses and component references. (gimplify_adjust_omp_clauses_1): Skip adjustments for explicit attach/detach clauses. (gimplify_omp_target_update): Handle finalize for detach. * omp-low.c (lower_omp_target): Support GOMP_MAP_ATTACH, GOMP_MAP_DETACH, GOMP_MAP_FORCE_DETACH. * tree-pretty-print.c (dump_omp_clause): Likewise, plus GOMP_MAP_ATTACH_DETACH. include/ * gomp-constants.h (GOMP_MAP_DEEP_COPY): Define. (gomp_map_kind): Add GOMP_MAP_ATTACH, GOMP_MAP_DETACH, GOMP_MAP_FORCE_DETACH. gcc/testsuite/ * c-c++-common/goacc/deep-copy-arrayofstruct.c: New test. * c-c++-common/goacc/mdc-1.c: New test. * c-c++-common/goacc/mdc-2.c: New test. * gcc.dg/goacc/mdc.C: New test. * gfortran.dg/goacc/data-clauses.f95: New test. * gfortran.dg/goacc/derived-types.f90: New test. * gfortran.dg/goacc/derived-types-2.f90: New test. * gfortran.dg/goacc/enter-exit-data.f95: New test. libgomp/ * libgomp.h (struct target_var_desc): Add do_detach flag. (struct splay_tree_key_s): Add attach_count field to union. (gomp_attach_pointer, gomp_detach_pointer): Add prototypes. * libgomp.map (OACC_2.6): New section. Add acc_attach, acc_attach_async, acc_detach, acc_detach_async, acc_detach_finalize, acc_detach_finalize_async. * oacc-mem.c (acc_attach_async, acc_attach, goacc_detach_internal, acc_detach, acc_detach_async, acc_detach_finalize, acc_detach_finalize_async): New functions. * oacc-parallel.c (find_pointer): Support attach/detach. (GOACC_enter_exit_data): Support attach/detach and GOMP_MAP_STRUCT. Don't call gomp_acc_insert_pointer. * openacc.h (acc_attach, acc_attach_async, acc_detach, (acc_detach_async, acc_detach_finalize, acc_detach_finalize_async): Add prototypes. * target.c (dump_tgt): Support attach_count field. (gomp_map_vars_existing): Initialise do_detach field of tgt_var_desc. (gomp_attach_pointer, gomp_detach_pointer): New functions. (gomp_map_vars_internal): Support attach and detach. (gomp_remove_var_internal): Free attach count array if present. (gomp_unmap_vars_internal): Support detach. (gomp_load_image_to_device): Zero-initialise attach_count field. * testsuite/libgomp.oacc-c-c++-common/deep-copy-1.c: New test. * testsuite/libgomp.oacc-c-c++-common/deep-copy-2.c: New test. * testsuite/libgomp.oacc-c-c++-common/deep-copy-3.c: New test. * testsuite/libgomp.oacc-c-c++-common/deep-copy-4.c: New test. * testsuite/libgomp.oacc-c-c++-common/deep-copy-5.c: New test. * testsuite/libgomp.oacc-c-c++-common/deep-copy-6.c: New test. * testsuite/libgomp.oacc-c-c++-common/deep-copy-7.c: New test. * testsuite/libgomp.oacc-c-c++-common/deep-copy-8.c: New test. * testsuite/libgomp.oacc-c-c++-common/deep-copy-9.c: New test. * testsuite/libgomp.oacc-c-c++-common/deep-copy-10.c: New test. * testsuite/libgomp.oacc-c-c++-common/deep-copy-11.c: New test. * testsuite/libgomp.oacc-c-c++-common/deep-copy-14.c: New test. * testsuite/libgomp.oacc-c++/deep-copy-12.C: Mew test. * testsuite/libgomp.oacc-c++/deep-copy-13.C: Mew test. * testsuite/libgomp.oacc-fortran/deep-copy-1.f90: New test. * testsuite/libgomp.oacc-fortran/deep-copy-2.f90: New test. * testsuite/libgomp.oacc-fortran/deep-copy-3.f90: New test. * testsuite/libgomp.oacc-fortran/deep-copy-4.f90: New test. * testsuite/libgomp.oacc-fortran/deep-copy-5.f90: New test. * testsuite/libgomp.oacc-fortran/deep-copy-6.f90: New test. * testsuite/libgomp.oacc-fortran/deep-copy-7.f90: New test. * testsuite/libgomp.oacc-fortran/deep-copy-8.f90: New test. * testsuite/libgomp.oacc-fortran/derived-type-1.f90: New test. * testsuite/libgomp.oacc-fortran/derivedtype-1.f95: New test. * testsuite/libgomp.oacc-fortran/derivedtype-2.f95: New test. * testsuite/libgomp.oacc-fortran/multidim-slice.f95: New test. * testsuite/libgomp.oacc-fortran/update-2.f90: New test. --- gcc/c-family/c-pragma.h | 2 + gcc/c/c-parser.c | 52 +++- gcc/c/c-typeck.c | 76 ++++- gcc/cp/parser.c | 55 +++- gcc/cp/semantics.c | 98 +++++- gcc/fortran/gfortran.h | 2 + gcc/fortran/openmp.c | 150 ++++++--- gcc/fortran/trans-expr.c | 184 +++++------ gcc/fortran/trans-openmp.c | 290 ++++++++++++++---- gcc/fortran/trans.h | 9 + gcc/gimplify.c | 224 ++++++++++++-- gcc/omp-low.c | 3 + .../goacc/deep-copy-arrayofstruct.c | 84 +++++ gcc/testsuite/c-c++-common/goacc/mdc-1.c | 55 ++++ gcc/testsuite/c-c++-common/goacc/mdc-2.c | 62 ++++ gcc/testsuite/g++.dg/goacc/mdc.C | 68 ++++ .../gfortran.dg/goacc/data-clauses.f95 | 38 +-- .../gfortran.dg/goacc/derived-types-2.f90 | 14 + .../gfortran.dg/goacc/derived-types.f90 | 77 +++++ .../gfortran.dg/goacc/enter-exit-data.f95 | 24 +- gcc/tree-pretty-print.c | 18 ++ include/gomp-constants.h | 16 +- libgomp/libgomp.h | 12 + libgomp/libgomp.map | 10 + libgomp/oacc-mem.c | 82 ++++- libgomp/oacc-parallel.c | 193 +++++++++--- libgomp/openacc.h | 6 + libgomp/target.c | 192 +++++++++++- .../testsuite/libgomp.oacc-c++/deep-copy-12.C | 72 +++++ .../testsuite/libgomp.oacc-c++/deep-copy-13.C | 72 +++++ .../libgomp.oacc-c-c++-common/deep-copy-1.c | 24 ++ .../libgomp.oacc-c-c++-common/deep-copy-10.c | 53 ++++ .../libgomp.oacc-c-c++-common/deep-copy-11.c | 72 +++++ .../libgomp.oacc-c-c++-common/deep-copy-14.c | 63 ++++ .../libgomp.oacc-c-c++-common/deep-copy-2.c | 29 ++ .../libgomp.oacc-c-c++-common/deep-copy-3.c | 34 ++ .../libgomp.oacc-c-c++-common/deep-copy-4.c | 87 ++++++ .../libgomp.oacc-c-c++-common/deep-copy-5.c | 81 +++++ .../libgomp.oacc-c-c++-common/deep-copy-6.c | 59 ++++ .../libgomp.oacc-c-c++-common/deep-copy-7.c | 45 +++ .../libgomp.oacc-c-c++-common/deep-copy-8.c | 54 ++++ .../libgomp.oacc-c-c++-common/deep-copy-9.c | 53 ++++ .../libgomp.oacc-fortran/deep-copy-1.f90 | 35 +++ .../libgomp.oacc-fortran/deep-copy-2.f90 | 33 ++ .../libgomp.oacc-fortran/deep-copy-3.f90 | 34 ++ .../libgomp.oacc-fortran/deep-copy-4.f90 | 49 +++ .../libgomp.oacc-fortran/deep-copy-5.f90 | 57 ++++ .../libgomp.oacc-fortran/deep-copy-6.f90 | 61 ++++ .../libgomp.oacc-fortran/deep-copy-7.f90 | 89 ++++++ .../libgomp.oacc-fortran/deep-copy-8.f90 | 41 +++ .../libgomp.oacc-fortran/derived-type-1.f90 | 28 ++ .../libgomp.oacc-fortran/derivedtype-1.f95 | 30 ++ .../libgomp.oacc-fortran/derivedtype-2.f95 | 41 +++ .../libgomp.oacc-fortran/multidim-slice.f95 | 50 +++ .../libgomp.oacc-fortran/update-2.f90 | 284 +++++++++++++++++ 55 files changed, 3389 insertions(+), 337 deletions(-) create mode 100644 gcc/testsuite/c-c++-common/goacc/deep-copy-arrayofstruct.c create mode 100644 gcc/testsuite/c-c++-common/goacc/mdc-1.c create mode 100644 gcc/testsuite/c-c++-common/goacc/mdc-2.c create mode 100644 gcc/testsuite/g++.dg/goacc/mdc.C create mode 100644 gcc/testsuite/gfortran.dg/goacc/derived-types-2.f90 create mode 100644 gcc/testsuite/gfortran.dg/goacc/derived-types.f90 create mode 100644 libgomp/testsuite/libgomp.oacc-c++/deep-copy-12.C create mode 100644 libgomp/testsuite/libgomp.oacc-c++/deep-copy-13.C create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-1.c create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-10.c create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-11.c create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-14.c create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-2.c create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-3.c create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-4.c create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-5.c create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-6.c create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-7.c create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-8.c create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-9.c create mode 100644 libgomp/testsuite/libgomp.oacc-fortran/deep-copy-1.f90 create mode 100644 libgomp/testsuite/libgomp.oacc-fortran/deep-copy-2.f90 create mode 100644 libgomp/testsuite/libgomp.oacc-fortran/deep-copy-3.f90 create mode 100644 libgomp/testsuite/libgomp.oacc-fortran/deep-copy-4.f90 create mode 100644 libgomp/testsuite/libgomp.oacc-fortran/deep-copy-5.f90 create mode 100644 libgomp/testsuite/libgomp.oacc-fortran/deep-copy-6.f90 create mode 100644 libgomp/testsuite/libgomp.oacc-fortran/deep-copy-7.f90 create mode 100644 libgomp/testsuite/libgomp.oacc-fortran/deep-copy-8.f90 create mode 100644 libgomp/testsuite/libgomp.oacc-fortran/derived-type-1.f90 create mode 100644 libgomp/testsuite/libgomp.oacc-fortran/derivedtype-1.f95 create mode 100644 libgomp/testsuite/libgomp.oacc-fortran/derivedtype-2.f95 create mode 100644 libgomp/testsuite/libgomp.oacc-fortran/multidim-slice.f95 create mode 100644 libgomp/testsuite/libgomp.oacc-fortran/update-2.f90 diff --git a/gcc/c-family/c-pragma.h b/gcc/c-family/c-pragma.h index e0aa774555a..0bbfd35d261 100644 --- a/gcc/c-family/c-pragma.h +++ b/gcc/c-family/c-pragma.h @@ -142,11 +142,13 @@ enum pragma_omp_clause { /* Clauses for OpenACC. */ PRAGMA_OACC_CLAUSE_ASYNC, + PRAGMA_OACC_CLAUSE_ATTACH, PRAGMA_OACC_CLAUSE_AUTO, PRAGMA_OACC_CLAUSE_COPY, PRAGMA_OACC_CLAUSE_COPYOUT, PRAGMA_OACC_CLAUSE_CREATE, PRAGMA_OACC_CLAUSE_DELETE, + PRAGMA_OACC_CLAUSE_DETACH, PRAGMA_OACC_CLAUSE_DEVICEPTR, PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT, PRAGMA_OACC_CLAUSE_FINALIZE, diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c index 6957297b3a5..85001394317 100644 --- a/gcc/c/c-parser.c +++ b/gcc/c/c-parser.c @@ -11745,6 +11745,8 @@ c_parser_omp_clause_name (c_parser *parser) result = PRAGMA_OMP_CLAUSE_ALIGNED; else if (!strcmp ("async", p)) result = PRAGMA_OACC_CLAUSE_ASYNC; + else if (!strcmp ("attach", p)) + result = PRAGMA_OACC_CLAUSE_ATTACH; break; case 'b': if (!strcmp ("bind", p)) @@ -11771,6 +11773,8 @@ c_parser_omp_clause_name (c_parser *parser) result = PRAGMA_OACC_CLAUSE_DELETE; else if (!strcmp ("depend", p)) result = PRAGMA_OMP_CLAUSE_DEPEND; + else if (!strcmp ("detach", p)) + result = PRAGMA_OACC_CLAUSE_DETACH; else if (!strcmp ("device", p)) result = PRAGMA_OMP_CLAUSE_DEVICE; else if (!strcmp ("deviceptr", p)) @@ -12014,12 +12018,16 @@ c_parser_oacc_wait_list (c_parser *parser, location_t clause_loc, tree list) If KIND is nonzero, CLAUSE_LOC is the location of the clause. If KIND is zero, create a TREE_LIST with the decl in TREE_PURPOSE; - return the list created. */ + return the list created. + + The optional ALLOW_DEREF argument is true if list items can use the deref + (->) operator. */ static tree c_parser_omp_variable_list (c_parser *parser, location_t clause_loc, - enum omp_clause_code kind, tree list) + enum omp_clause_code kind, tree list, + bool allow_deref = false) { auto_vec tokens; unsigned int tokens_avail = 0; @@ -12146,9 +12154,13 @@ c_parser_omp_variable_list (c_parser *parser, case OMP_CLAUSE_MAP: case OMP_CLAUSE_FROM: case OMP_CLAUSE_TO: - while (c_parser_next_token_is (parser, CPP_DOT)) + while (c_parser_next_token_is (parser, CPP_DOT) + || (allow_deref + && c_parser_next_token_is (parser, CPP_DEREF))) { location_t op_loc = c_parser_peek_token (parser)->location; + if (c_parser_next_token_is (parser, CPP_DEREF)) + t = build_simple_mem_ref (t); c_parser_consume_token (parser); if (!c_parser_next_token_is (parser, CPP_NAME)) { @@ -12270,11 +12282,12 @@ c_parser_omp_variable_list (c_parser *parser, } /* Similarly, but expect leading and trailing parenthesis. This is a very - common case for OpenACC and OpenMP clauses. */ + common case for OpenACC and OpenMP clauses. The optional ALLOW_DEREF + argument is true if list items can use the deref (->) operator. */ static tree c_parser_omp_var_list_parens (c_parser *parser, enum omp_clause_code kind, - tree list) + tree list, bool allow_deref = false) { /* The clauses location. */ location_t loc = c_parser_peek_token (parser)->location; @@ -12282,18 +12295,20 @@ c_parser_omp_var_list_parens (c_parser *parser, enum omp_clause_code kind, matching_parens parens; if (parens.require_open (parser)) { - list = c_parser_omp_variable_list (parser, loc, kind, list); + list = c_parser_omp_variable_list (parser, loc, kind, list, allow_deref); parens.skip_until_found_close (parser); } return list; } -/* OpenACC 2.0: +/* OpenACC 2.0+: + attach ( variable-list ) copy ( variable-list ) copyin ( variable-list ) copyout ( variable-list ) create ( variable-list ) delete ( variable-list ) + detach ( variable-list ) present ( variable-list ) */ static tree @@ -12303,6 +12318,9 @@ c_parser_oacc_data_clause (c_parser *parser, pragma_omp_clause c_kind, enum gomp_map_kind kind; switch (c_kind) { + case PRAGMA_OACC_CLAUSE_ATTACH: + kind = GOMP_MAP_ATTACH; + break; case PRAGMA_OACC_CLAUSE_COPY: kind = GOMP_MAP_TOFROM; break; @@ -12318,6 +12336,9 @@ c_parser_oacc_data_clause (c_parser *parser, pragma_omp_clause c_kind, case PRAGMA_OACC_CLAUSE_DELETE: kind = GOMP_MAP_RELEASE; break; + case PRAGMA_OACC_CLAUSE_DETACH: + kind = GOMP_MAP_DETACH; + break; case PRAGMA_OACC_CLAUSE_DEVICE: kind = GOMP_MAP_FORCE_TO; break; @@ -12337,7 +12358,7 @@ c_parser_oacc_data_clause (c_parser *parser, pragma_omp_clause c_kind, gcc_unreachable (); } tree nl, c; - nl = c_parser_omp_var_list_parens (parser, OMP_CLAUSE_MAP, list); + nl = c_parser_omp_var_list_parens (parser, OMP_CLAUSE_MAP, list, true); for (c = nl; c != list; c = OMP_CLAUSE_CHAIN (c)) OMP_CLAUSE_SET_MAP_KIND (c, kind); @@ -15052,6 +15073,10 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask, clauses); c_name = "auto"; break; + case PRAGMA_OACC_CLAUSE_ATTACH: + clauses = c_parser_oacc_data_clause (parser, c_kind, clauses); + c_name = "attach"; + break; case PRAGMA_OACC_CLAUSE_COLLAPSE: clauses = c_parser_omp_clause_collapse (parser, clauses); c_name = "collapse"; @@ -15080,6 +15105,10 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask, clauses = c_parser_omp_clause_default (parser, clauses, true); c_name = "default"; break; + case PRAGMA_OACC_CLAUSE_DETACH: + clauses = c_parser_oacc_data_clause (parser, c_kind, clauses); + c_name = "detach"; + break; case PRAGMA_OACC_CLAUSE_DEVICE: clauses = c_parser_oacc_data_clause (parser, c_kind, clauses); c_name = "device"; @@ -15574,7 +15603,8 @@ c_parser_oacc_cache (location_t loc, c_parser *parser) */ #define OACC_DATA_CLAUSE_MASK \ - ( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPY) \ + ( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ATTACH) \ + | (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) \ @@ -15757,6 +15787,7 @@ c_parser_oacc_declare (c_parser *parser) #define OACC_ENTER_DATA_CLAUSE_MASK \ ( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ASYNC) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ATTACH) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYIN) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_CREATE) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) ) @@ -15766,6 +15797,7 @@ c_parser_oacc_declare (c_parser *parser) | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ASYNC) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYOUT) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DELETE) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DETACH) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_FINALIZE) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) ) @@ -15900,6 +15932,7 @@ c_parser_oacc_loop (location_t loc, c_parser *parser, char *p_name, #define OACC_KERNELS_CLAUSE_MASK \ ( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ASYNC) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ATTACH) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPY) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYIN) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYOUT) \ @@ -15915,6 +15948,7 @@ c_parser_oacc_loop (location_t loc, c_parser *parser, char *p_name, #define OACC_PARALLEL_CLAUSE_MASK \ ( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ASYNC) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ATTACH) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPY) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYIN) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYOUT) \ diff --git a/gcc/c/c-typeck.c b/gcc/c/c-typeck.c index c7339509bd1..0b9b207e653 100644 --- a/gcc/c/c-typeck.c +++ b/gcc/c/c-typeck.c @@ -12884,7 +12884,6 @@ handle_omp_array_sections_1 (tree c, tree t, vec &types, return error_mark_node; } if (TREE_CODE (t) == COMPONENT_REF - && ort == C_ORT_OMP && (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_TO || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FROM)) @@ -12905,6 +12904,15 @@ handle_omp_array_sections_1 (tree c, tree t, vec &types, return error_mark_node; } t = TREE_OPERAND (t, 0); + if (ort == C_ORT_ACC && TREE_CODE (t) == MEM_REF) + { + if (maybe_ne (mem_ref_offset (t), 0)) + error_at (OMP_CLAUSE_LOCATION (c), + "cannot dereference %qE in %qs clause", t, + omp_clause_code_name[OMP_CLAUSE_CODE (c)]); + else + t = TREE_OPERAND (t, 0); + } } } if (!VAR_P (t) && TREE_CODE (t) != PARM_DECL) @@ -12990,7 +12998,18 @@ handle_omp_array_sections_1 (tree c, tree t, vec &types, length = fold_convert (sizetype, length); if (low_bound == NULL_TREE) low_bound = integer_zero_node; - + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP + && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH + || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH)) + { + if (length != integer_one_node) + { + error_at (OMP_CLAUSE_LOCATION (c), + "expected single pointer in %qs clause", + c_omp_map_clause_name (c, ort == C_ORT_ACC)); + return error_mark_node; + } + } if (length != NULL_TREE) { if (!integer_nonzerop (length)) @@ -13430,7 +13449,11 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort) if (ort != C_ORT_OMP && ort != C_ORT_ACC) OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_POINTER); else if (TREE_CODE (t) == COMPONENT_REF) - OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ALWAYS_POINTER); + { + gomp_map_kind k = (ort == C_ORT_ACC) ? GOMP_MAP_ATTACH_DETACH + : GOMP_MAP_ALWAYS_POINTER; + OMP_CLAUSE_SET_MAP_KIND (c2, k); + } else OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_FIRSTPRIVATE_POINTER); if (OMP_CLAUSE_MAP_KIND (c2) != GOMP_MAP_FIRSTPRIVATE_POINTER @@ -13667,6 +13690,35 @@ c_omp_finish_iterators (tree iter) return ret; } +/* Ensure that pointers are used in OpenACC attach and detach clauses. + Return true if an error has been detected. */ + +static bool +c_oacc_check_attachments (tree c) +{ + if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP) + return false; + + /* OpenACC attach / detach clauses must be pointers. */ + if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH + || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH) + { + tree t = OMP_CLAUSE_DECL (c); + + while (TREE_CODE (t) == TREE_LIST) + t = TREE_CHAIN (t); + + if (TREE_CODE (TREE_TYPE (t)) != POINTER_TYPE) + { + error_at (OMP_CLAUSE_LOCATION (c), "expected pointer in %qs clause", + c_omp_map_clause_name (c, true)); + return true; + } + } + + return false; +} + /* For all elements of CLAUSES, validate them against their constraints. Remove any elements from the list that are invalid. */ @@ -14420,6 +14472,8 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) } } } + if (c_oacc_check_attachments (c)) + remove = true; break; } if (t == error_mark_node) @@ -14427,8 +14481,13 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) remove = true; break; } + /* OpenACC attach / detach clauses must be pointers. */ + if (c_oacc_check_attachments (c)) + { + remove = true; + break; + } if (TREE_CODE (t) == COMPONENT_REF - && (ort & C_ORT_OMP) && OMP_CLAUSE_CODE (c) != OMP_CLAUSE__CACHE_) { if (DECL_BIT_FIELD (TREE_OPERAND (t, 1))) @@ -14463,6 +14522,15 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) break; } t = TREE_OPERAND (t, 0); + if (ort == C_ORT_ACC && TREE_CODE (t) == MEM_REF) + { + if (maybe_ne (mem_ref_offset (t), 0)) + error_at (OMP_CLAUSE_LOCATION (c), + "cannot dereference %qE in %qs clause", t, + omp_clause_code_name[OMP_CLAUSE_CODE (c)]); + else + t = TREE_OPERAND (t, 0); + } } if (remove) break; diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c index 018a03d4965..d3e71e38983 100644 --- a/gcc/cp/parser.c +++ b/gcc/cp/parser.c @@ -32480,6 +32480,8 @@ cp_parser_omp_clause_name (cp_parser *parser) result = PRAGMA_OMP_CLAUSE_ALIGNED; else if (!strcmp ("async", p)) result = PRAGMA_OACC_CLAUSE_ASYNC; + else if (!strcmp ("attach", p)) + result = PRAGMA_OACC_CLAUSE_ATTACH; break; case 'b': if (!strcmp ("bind", p)) @@ -32504,6 +32506,8 @@ cp_parser_omp_clause_name (cp_parser *parser) result = PRAGMA_OMP_CLAUSE_DEFAULTMAP; else if (!strcmp ("depend", p)) result = PRAGMA_OMP_CLAUSE_DEPEND; + else if (!strcmp ("detach", p)) + result = PRAGMA_OACC_CLAUSE_DETACH; else if (!strcmp ("device", p)) result = PRAGMA_OMP_CLAUSE_DEVICE; else if (!strcmp ("deviceptr", p)) @@ -32706,11 +32710,15 @@ check_no_duplicate_clause (tree clauses, enum omp_clause_code code, COLON can be NULL if only closing parenthesis should end the list, or pointer to bool which will receive false if the list is terminated - by closing parenthesis or true if the list is terminated by colon. */ + by closing parenthesis or true if the list is terminated by colon. + + The optional ALLOW_DEREF argument is true if list items can use the deref + (->) operator. */ static tree cp_parser_omp_var_list_no_open (cp_parser *parser, enum omp_clause_code kind, - tree list, bool *colon) + tree list, bool *colon, + bool allow_deref = false) { cp_token *token; bool saved_colon_corrects_to_scope_p = parser->colon_corrects_to_scope_p; @@ -32791,15 +32799,20 @@ cp_parser_omp_var_list_no_open (cp_parser *parser, enum omp_clause_code kind, case OMP_CLAUSE_MAP: case OMP_CLAUSE_FROM: case OMP_CLAUSE_TO: - while (cp_lexer_next_token_is (parser->lexer, CPP_DOT)) + while (cp_lexer_next_token_is (parser->lexer, CPP_DOT) + || (allow_deref + && cp_lexer_next_token_is (parser->lexer, CPP_DEREF))) { + cpp_ttype ttype + = cp_lexer_next_token_is (parser->lexer, CPP_DOT) + ? CPP_DOT : CPP_DEREF; location_t loc = cp_lexer_peek_token (parser->lexer)->location; cp_id_kind idk = CP_ID_KIND_NONE; cp_lexer_consume_token (parser->lexer); decl = convert_from_reference (decl); decl - = cp_parser_postfix_dot_deref_expression (parser, CPP_DOT, + = cp_parser_postfix_dot_deref_expression (parser, ttype, decl, false, &idk, loc); } @@ -32917,19 +32930,23 @@ cp_parser_omp_var_list_no_open (cp_parser *parser, enum omp_clause_code kind, common case for omp clauses. */ static tree -cp_parser_omp_var_list (cp_parser *parser, enum omp_clause_code kind, tree list) +cp_parser_omp_var_list (cp_parser *parser, enum omp_clause_code kind, tree list, + bool allow_deref = false) { if (cp_parser_require (parser, CPP_OPEN_PAREN, RT_OPEN_PAREN)) - return cp_parser_omp_var_list_no_open (parser, kind, list, NULL); + return cp_parser_omp_var_list_no_open (parser, kind, list, NULL, + allow_deref); return list; } -/* OpenACC 2.0: +/* OpenACC 2.0+: + attach ( variable-list ) copy ( variable-list ) copyin ( variable-list ) copyout ( variable-list ) create ( variable-list ) delete ( variable-list ) + detach ( variable-list ) present ( variable-list ) */ static tree @@ -32939,6 +32956,9 @@ cp_parser_oacc_data_clause (cp_parser *parser, pragma_omp_clause c_kind, enum gomp_map_kind kind; switch (c_kind) { + case PRAGMA_OACC_CLAUSE_ATTACH: + kind = GOMP_MAP_ATTACH; + break; case PRAGMA_OACC_CLAUSE_COPY: kind = GOMP_MAP_TOFROM; break; @@ -32954,6 +32974,9 @@ cp_parser_oacc_data_clause (cp_parser *parser, pragma_omp_clause c_kind, case PRAGMA_OACC_CLAUSE_DELETE: kind = GOMP_MAP_RELEASE; break; + case PRAGMA_OACC_CLAUSE_DETACH: + kind = GOMP_MAP_DETACH; + break; case PRAGMA_OACC_CLAUSE_DEVICE: kind = GOMP_MAP_FORCE_TO; break; @@ -32973,7 +32996,7 @@ cp_parser_oacc_data_clause (cp_parser *parser, pragma_omp_clause c_kind, gcc_unreachable (); } tree nl, c; - nl = cp_parser_omp_var_list (parser, OMP_CLAUSE_MAP, list); + nl = cp_parser_omp_var_list (parser, OMP_CLAUSE_MAP, list, true); for (c = nl; c != list; c = OMP_CLAUSE_CHAIN (c)) OMP_CLAUSE_SET_MAP_KIND (c, kind); @@ -35451,6 +35474,10 @@ cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask, clauses); c_name = "auto"; break; + case PRAGMA_OACC_CLAUSE_ATTACH: + clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses); + c_name = "attach"; + break; case PRAGMA_OACC_CLAUSE_COLLAPSE: clauses = cp_parser_omp_clause_collapse (parser, clauses, here); c_name = "collapse"; @@ -35479,6 +35506,10 @@ cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask, clauses = cp_parser_omp_clause_default (parser, clauses, here, true); c_name = "default"; break; + case PRAGMA_OACC_CLAUSE_DETACH: + clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses); + c_name = "detach"; + break; case PRAGMA_OACC_CLAUSE_DEVICE: clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses); c_name = "device"; @@ -39313,10 +39344,12 @@ cp_parser_oacc_cache (cp_parser *parser, cp_token *pragma_tok) structured-block */ #define OACC_DATA_CLAUSE_MASK \ - ( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPY) \ + ( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ATTACH) \ + | (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_DETACH) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICEPTR) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT) ) @@ -39516,6 +39549,7 @@ cp_parser_oacc_declare (cp_parser *parser, cp_token *pragma_tok) #define OACC_ENTER_DATA_CLAUSE_MASK \ ( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ATTACH) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ASYNC) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYIN) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_CREATE) \ @@ -39526,6 +39560,7 @@ cp_parser_oacc_declare (cp_parser *parser, cp_token *pragma_tok) | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ASYNC) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYOUT) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DELETE) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DETACH) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_FINALIZE) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) ) @@ -39629,6 +39664,7 @@ cp_parser_oacc_loop (cp_parser *parser, cp_token *pragma_tok, char *p_name, #define OACC_KERNELS_CLAUSE_MASK \ ( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ASYNC) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ATTACH) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPY) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYIN) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYOUT) \ @@ -39644,6 +39680,7 @@ cp_parser_oacc_loop (cp_parser *parser, cp_token *pragma_tok, char *p_name, #define OACC_PARALLEL_CLAUSE_MASK \ ( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ASYNC) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ATTACH) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPY) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYIN) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYOUT) \ diff --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c index 5a8c4d237e0..f48bff42fc2 100644 --- a/gcc/cp/semantics.c +++ b/gcc/cp/semantics.c @@ -4700,7 +4700,6 @@ handle_omp_array_sections_1 (tree c, tree t, vec &types, t = TREE_OPERAND (t, 0); ret = t; if (TREE_CODE (t) == COMPONENT_REF - && ort == C_ORT_OMP && (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_TO || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FROM) @@ -4724,6 +4723,8 @@ handle_omp_array_sections_1 (tree c, tree t, vec &types, return error_mark_node; } t = TREE_OPERAND (t, 0); + if (ort == C_ORT_ACC && TREE_CODE (t) == INDIRECT_REF) + t = TREE_OPERAND (t, 0); } if (REFERENCE_REF_P (t)) t = TREE_OPERAND (t, 0); @@ -4823,6 +4824,18 @@ handle_omp_array_sections_1 (tree c, tree t, vec &types, if (low_bound == NULL_TREE) low_bound = integer_zero_node; + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP + && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH + || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH)) + { + if (length != integer_one_node) + { + error_at (OMP_CLAUSE_LOCATION (c), + "expected single pointer in %qs clause", + c_omp_map_clause_name (c, ort == C_ORT_ACC)); + return error_mark_node; + } + } if (length != NULL_TREE) { if (!integer_nonzerop (length)) @@ -5259,12 +5272,18 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort) if ((ort & C_ORT_OMP_DECLARE_SIMD) != C_ORT_OMP && ort != C_ORT_ACC) OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_POINTER); else if (TREE_CODE (t) == COMPONENT_REF) - OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ALWAYS_POINTER); + { + gomp_map_kind k = (ort == C_ORT_ACC) ? GOMP_MAP_ATTACH_DETACH + : GOMP_MAP_ALWAYS_POINTER; + OMP_CLAUSE_SET_MAP_KIND (c2, k); + } else if (REFERENCE_REF_P (t) && TREE_CODE (TREE_OPERAND (t, 0)) == COMPONENT_REF) { t = TREE_OPERAND (t, 0); - OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ALWAYS_POINTER); + gomp_map_kind k = (ort == C_ORT_ACC) ? GOMP_MAP_ATTACH_DETACH + : GOMP_MAP_ALWAYS_POINTER; + OMP_CLAUSE_SET_MAP_KIND (c2, k); } else OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_FIRSTPRIVATE_POINTER); @@ -6180,6 +6199,41 @@ cp_omp_finish_iterators (tree iter) return ret; } +/* Ensure that pointers are used in OpenACC attach and detach clauses. + Return true if an error has been detected. */ + +static bool +cp_oacc_check_attachments (tree c) +{ + if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP) + return false; + + /* OpenACC attach / detach clauses must be pointers. */ + if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH + || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH) + { + tree t = OMP_CLAUSE_DECL (c); + tree type; + + while (TREE_CODE (t) == TREE_LIST) + t = TREE_CHAIN (t); + + type = TREE_TYPE (t); + + if (TREE_CODE (type) == REFERENCE_TYPE) + type = TREE_TYPE (type); + + if (TREE_CODE (type) != POINTER_TYPE) + { + error_at (OMP_CLAUSE_LOCATION (c), "expected pointer in %qs clause", + c_omp_map_clause_name (c, true)); + return true; + } + } + + return false; +} + /* For all elements of CLAUSES, validate them vs OpenMP constraints. Remove any elements from the list that are invalid. */ @@ -6444,7 +6498,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) t = OMP_CLAUSE_DECL (c); check_dup_generic_t: if (t == current_class_ptr - && (ort != C_ORT_OMP_DECLARE_SIMD + && ((ort != C_ORT_OMP_DECLARE_SIMD && ort != C_ORT_ACC) || (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_LINEAR && OMP_CLAUSE_CODE (c) != OMP_CLAUSE_UNIFORM))) { @@ -6514,8 +6568,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) handle_field_decl: if (!remove && TREE_CODE (t) == FIELD_DECL - && t == OMP_CLAUSE_DECL (c) - && ort != C_ORT_ACC) + && t == OMP_CLAUSE_DECL (c)) { OMP_CLAUSE_DECL (c) = omp_privatize_field (t, (OMP_CLAUSE_CODE (c) @@ -6582,7 +6635,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) omp_note_field_privatization (t, OMP_CLAUSE_DECL (c)); else t = OMP_CLAUSE_DECL (c); - if (t == current_class_ptr) + if (ort != C_ORT_ACC && t == current_class_ptr) { error_at (OMP_CLAUSE_LOCATION (c), "% allowed in OpenMP only in %" @@ -7071,7 +7124,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) } if (t == error_mark_node) remove = true; - else if (t == current_class_ptr) + else if (ort != C_ORT_ACC && t == current_class_ptr) { error_at (OMP_CLAUSE_LOCATION (c), "% allowed in OpenMP only in %" @@ -7202,6 +7255,8 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) } } } + if (cp_oacc_check_attachments (c)) + remove = true; break; } if (t == error_mark_node) @@ -7209,14 +7264,25 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) remove = true; break; } + /* OpenACC attach / detach clauses must be pointers. */ + if (cp_oacc_check_attachments (c)) + { + remove = true; + break; + } if (REFERENCE_REF_P (t) && TREE_CODE (TREE_OPERAND (t, 0)) == COMPONENT_REF) { t = TREE_OPERAND (t, 0); OMP_CLAUSE_DECL (c) = t; } + if (ort == C_ORT_ACC + && TREE_CODE (t) == COMPONENT_REF + && TREE_CODE (TREE_OPERAND (t, 0)) == INDIRECT_REF) + t = TREE_OPERAND (TREE_OPERAND (t, 0), 0); if (TREE_CODE (t) == COMPONENT_REF - && (ort & C_ORT_OMP_DECLARE_SIMD) == C_ORT_OMP + && ((ort & C_ORT_OMP_DECLARE_SIMD) == C_ORT_OMP + || ort == C_ORT_ACC) && OMP_CLAUSE_CODE (c) != OMP_CLAUSE__CACHE_) { if (type_dependent_expression_p (t)) @@ -7266,7 +7332,8 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) break; if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER - || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_POINTER)) + || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_POINTER + || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH)) break; if (DECL_P (t)) error_at (OMP_CLAUSE_LOCATION (c), @@ -7348,7 +7415,9 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) else bitmap_set_bit (&generic_head, DECL_UID (t)); } - else if (bitmap_bit_p (&map_head, DECL_UID (t))) + else if (bitmap_bit_p (&map_head, DECL_UID (t)) + && (ort != C_ORT_ACC + || !bitmap_bit_p (&map_field_head, DECL_UID (t)))) { if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP) error_at (OMP_CLAUSE_LOCATION (c), @@ -7403,7 +7472,12 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP); if (TREE_CODE (t) == COMPONENT_REF) - OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ALWAYS_POINTER); + { + gomp_map_kind k + = (ort == C_ORT_ACC) ? GOMP_MAP_ATTACH_DETACH + : GOMP_MAP_ALWAYS_POINTER; + OMP_CLAUSE_SET_MAP_KIND (c2, k); + } else OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_FIRSTPRIVATE_REFERENCE); diff --git a/gcc/fortran/gfortran.h b/gcc/fortran/gfortran.h index d84d1fa7f7e..45fdfdd4dd3 100644 --- a/gcc/fortran/gfortran.h +++ b/gcc/fortran/gfortran.h @@ -1191,10 +1191,12 @@ enum gfc_omp_depend_op enum gfc_omp_map_op { OMP_MAP_ALLOC, + OMP_MAP_ATTACH, OMP_MAP_TO, OMP_MAP_FROM, OMP_MAP_TOFROM, OMP_MAP_DELETE, + OMP_MAP_DETACH, OMP_MAP_FORCE_ALLOC, OMP_MAP_FORCE_TO, OMP_MAP_FORCE_FROM, diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c index 7df7384c187..f08f77ce940 100644 --- a/gcc/fortran/openmp.c +++ b/gcc/fortran/openmp.c @@ -222,7 +222,8 @@ static match gfc_match_omp_variable_list (const char *str, gfc_omp_namelist **list, bool allow_common, bool *end_colon = NULL, gfc_omp_namelist ***headp = NULL, - bool allow_sections = false) + bool allow_sections = false, + bool allow_derived = false) { gfc_omp_namelist *head, *tail, *p; locus old_loc, cur_loc; @@ -248,7 +249,8 @@ gfc_match_omp_variable_list (const char *str, gfc_omp_namelist **list, case MATCH_YES: gfc_expr *expr; expr = NULL; - if (allow_sections && gfc_peek_ascii_char () == '(') + if ((allow_sections && gfc_peek_ascii_char () == '(') + || (allow_derived && gfc_peek_ascii_char () == '%')) { gfc_current_locus = cur_loc; m = gfc_match_variable (&expr, 0); @@ -786,7 +788,7 @@ enum omp_mask1 OMP_MASK1_LAST }; -/* OpenACC 2.0 specific clauses. */ +/* OpenACC 2.0+ specific clauses. */ enum omp_mask2 { OMP_CLAUSE_ASYNC, @@ -812,6 +814,8 @@ enum omp_mask2 OMP_CLAUSE_TILE, OMP_CLAUSE_IF_PRESENT, OMP_CLAUSE_FINALIZE, + OMP_CLAUSE_ATTACH, + OMP_CLAUSE_DETACH, /* This must come last. */ OMP_MASK2_LAST }; @@ -915,10 +919,12 @@ omp_inv_mask::omp_inv_mask (const omp_mask &m) : omp_mask (m) mapping. */ static bool -gfc_match_omp_map_clause (gfc_omp_namelist **list, gfc_omp_map_op map_op) +gfc_match_omp_map_clause (gfc_omp_namelist **list, gfc_omp_map_op map_op, + bool allow_derived = false) { gfc_omp_namelist **head = NULL; - if (gfc_match_omp_variable_list ("", list, false, NULL, &head, true) + if (gfc_match_omp_variable_list ("", list, false, NULL, &head, true, + allow_derived) == MATCH_YES) { gfc_omp_namelist *n; @@ -940,6 +946,14 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask, { gfc_omp_clauses *c = gfc_get_omp_clauses (); locus old_loc; + /* Determine whether we're dealing with an OpenACC directive that permits + derived type member accesses. This in particular disallows + "!$acc declare" from using such accesses, because it's not clear if/how + that should work. */ + bool allow_derived = (openacc + && ((mask & OMP_CLAUSE_ATTACH) + || (mask & OMP_CLAUSE_DETACH) + || (mask & OMP_CLAUSE_HOST_SELF))); gcc_checking_assert (OMP_MASK1_LAST <= 64 && OMP_MASK2_LAST <= 64); *cp = NULL; @@ -1013,6 +1027,11 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask, needs_space = true; continue; } + if ((mask & OMP_CLAUSE_ATTACH) + && gfc_match ("attach ( ") == MATCH_YES + && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP], + OMP_MAP_ATTACH, allow_derived)) + continue; break; case 'c': if ((mask & OMP_CLAUSE_COLLAPSE) @@ -1040,7 +1059,7 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask, if ((mask & OMP_CLAUSE_COPY) && gfc_match ("copy ( ") == MATCH_YES && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP], - OMP_MAP_TOFROM)) + OMP_MAP_TOFROM, allow_derived)) continue; if (mask & OMP_CLAUSE_COPYIN) { @@ -1048,7 +1067,7 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask, { if (gfc_match ("copyin ( ") == MATCH_YES && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP], - OMP_MAP_TO)) + OMP_MAP_TO, allow_derived)) continue; } else if (gfc_match_omp_variable_list ("copyin (", @@ -1059,7 +1078,7 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask, if ((mask & OMP_CLAUSE_COPYOUT) && gfc_match ("copyout ( ") == MATCH_YES && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP], - OMP_MAP_FROM)) + OMP_MAP_FROM, allow_derived)) continue; if ((mask & OMP_CLAUSE_COPYPRIVATE) && gfc_match_omp_variable_list ("copyprivate (", @@ -1069,7 +1088,7 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask, if ((mask & OMP_CLAUSE_CREATE) && gfc_match ("create ( ") == MATCH_YES && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP], - OMP_MAP_ALLOC)) + OMP_MAP_ALLOC, allow_derived)) continue; break; case 'd': @@ -1105,7 +1124,7 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask, if ((mask & OMP_CLAUSE_DELETE) && gfc_match ("delete ( ") == MATCH_YES && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP], - OMP_MAP_RELEASE)) + OMP_MAP_RELEASE, allow_derived)) continue; if ((mask & OMP_CLAUSE_DEPEND) && gfc_match ("depend ( ") == MATCH_YES) @@ -1148,6 +1167,11 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask, else gfc_current_locus = old_loc; } + if ((mask & OMP_CLAUSE_DETACH) + && gfc_match ("detach ( ") == MATCH_YES + && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP], + OMP_MAP_DETACH, allow_derived)) + continue; if ((mask & OMP_CLAUSE_DEVICE) && !openacc && c->device == NULL @@ -1157,12 +1181,13 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask, && openacc && gfc_match ("device ( ") == MATCH_YES && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP], - OMP_MAP_FORCE_TO)) + OMP_MAP_FORCE_TO, allow_derived)) continue; if ((mask & OMP_CLAUSE_DEVICEPTR) && gfc_match ("deviceptr ( ") == MATCH_YES && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP], - OMP_MAP_FORCE_DEVICEPTR)) + OMP_MAP_FORCE_DEVICEPTR, + allow_derived)) continue; if ((mask & OMP_CLAUSE_DEVICE_RESIDENT) && gfc_match_omp_variable_list @@ -1240,7 +1265,7 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask, if ((mask & OMP_CLAUSE_HOST_SELF) && gfc_match ("host ( ") == MATCH_YES && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP], - OMP_MAP_FORCE_FROM)) + OMP_MAP_FORCE_FROM, allow_derived)) continue; break; case 'i': @@ -1512,47 +1537,48 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask, if ((mask & OMP_CLAUSE_COPY) && gfc_match ("pcopy ( ") == MATCH_YES && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP], - OMP_MAP_TOFROM)) + OMP_MAP_TOFROM, allow_derived)) continue; if ((mask & OMP_CLAUSE_COPYIN) && gfc_match ("pcopyin ( ") == MATCH_YES && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP], - OMP_MAP_TO)) + OMP_MAP_TO, allow_derived)) continue; if ((mask & OMP_CLAUSE_COPYOUT) && gfc_match ("pcopyout ( ") == MATCH_YES && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP], - OMP_MAP_FROM)) + OMP_MAP_FROM, allow_derived)) continue; if ((mask & OMP_CLAUSE_CREATE) && gfc_match ("pcreate ( ") == MATCH_YES && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP], - OMP_MAP_ALLOC)) + OMP_MAP_ALLOC, allow_derived)) continue; if ((mask & OMP_CLAUSE_PRESENT) && gfc_match ("present ( ") == MATCH_YES && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP], - OMP_MAP_FORCE_PRESENT)) + OMP_MAP_FORCE_PRESENT, + allow_derived)) continue; if ((mask & OMP_CLAUSE_COPY) && gfc_match ("present_or_copy ( ") == MATCH_YES && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP], - OMP_MAP_TOFROM)) + OMP_MAP_TOFROM, allow_derived)) continue; if ((mask & OMP_CLAUSE_COPYIN) && gfc_match ("present_or_copyin ( ") == MATCH_YES && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP], - OMP_MAP_TO)) + OMP_MAP_TO, allow_derived)) continue; if ((mask & OMP_CLAUSE_COPYOUT) && gfc_match ("present_or_copyout ( ") == MATCH_YES && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP], - OMP_MAP_FROM)) + OMP_MAP_FROM, allow_derived)) continue; if ((mask & OMP_CLAUSE_CREATE) && gfc_match ("present_or_create ( ") == MATCH_YES && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP], - OMP_MAP_ALLOC)) + OMP_MAP_ALLOC, allow_derived)) continue; if ((mask & OMP_CLAUSE_PRIORITY) && c->priority == NULL @@ -1670,8 +1696,8 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask, if (gfc_match_omp_variable_list (" :", &c->lists[OMP_LIST_REDUCTION], - false, NULL, &head, - openacc) == MATCH_YES) + false, NULL, &head, openacc, + allow_derived) == MATCH_YES) { gfc_omp_namelist *n; if (rop == OMP_REDUCTION_NONE) @@ -1770,7 +1796,7 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask, if ((mask & OMP_CLAUSE_HOST_SELF) && gfc_match ("self ( ") == MATCH_YES && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP], - OMP_MAP_FORCE_FROM)) + OMP_MAP_FORCE_FROM, allow_derived)) continue; if ((mask & OMP_CLAUSE_SEQ) && !c->seq @@ -1945,17 +1971,17 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask, | OMP_CLAUSE_COPY | OMP_CLAUSE_COPYIN | OMP_CLAUSE_COPYOUT \ | OMP_CLAUSE_CREATE | OMP_CLAUSE_PRESENT | OMP_CLAUSE_DEVICEPTR \ | OMP_CLAUSE_PRIVATE | OMP_CLAUSE_FIRSTPRIVATE | OMP_CLAUSE_DEFAULT \ - | OMP_CLAUSE_WAIT) + | OMP_CLAUSE_WAIT | OMP_CLAUSE_ATTACH) #define OACC_KERNELS_CLAUSES \ (omp_mask (OMP_CLAUSE_IF) | OMP_CLAUSE_ASYNC | OMP_CLAUSE_NUM_GANGS \ | OMP_CLAUSE_NUM_WORKERS | OMP_CLAUSE_VECTOR_LENGTH | OMP_CLAUSE_DEVICEPTR \ | OMP_CLAUSE_COPY | OMP_CLAUSE_COPYIN | OMP_CLAUSE_COPYOUT \ | OMP_CLAUSE_CREATE | OMP_CLAUSE_PRESENT | OMP_CLAUSE_DEFAULT \ - | OMP_CLAUSE_WAIT) + | OMP_CLAUSE_WAIT | OMP_CLAUSE_ATTACH) #define OACC_DATA_CLAUSES \ (omp_mask (OMP_CLAUSE_IF) | OMP_CLAUSE_DEVICEPTR | OMP_CLAUSE_COPY \ | OMP_CLAUSE_COPYIN | OMP_CLAUSE_COPYOUT | OMP_CLAUSE_CREATE \ - | OMP_CLAUSE_PRESENT) + | OMP_CLAUSE_PRESENT | OMP_CLAUSE_ATTACH) #define OACC_LOOP_CLAUSES \ (omp_mask (OMP_CLAUSE_COLLAPSE) | OMP_CLAUSE_GANG | OMP_CLAUSE_WORKER \ | OMP_CLAUSE_VECTOR | OMP_CLAUSE_SEQ | OMP_CLAUSE_INDEPENDENT \ @@ -1976,10 +2002,11 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask, | OMP_CLAUSE_DEVICE | OMP_CLAUSE_WAIT | OMP_CLAUSE_IF_PRESENT) #define OACC_ENTER_DATA_CLAUSES \ (omp_mask (OMP_CLAUSE_IF) | OMP_CLAUSE_ASYNC | OMP_CLAUSE_WAIT \ - | OMP_CLAUSE_COPYIN | OMP_CLAUSE_CREATE) + | OMP_CLAUSE_COPYIN | OMP_CLAUSE_CREATE | OMP_CLAUSE_ATTACH) #define OACC_EXIT_DATA_CLAUSES \ (omp_mask (OMP_CLAUSE_IF) | OMP_CLAUSE_ASYNC | OMP_CLAUSE_WAIT \ - | OMP_CLAUSE_COPYOUT | OMP_CLAUSE_DELETE | OMP_CLAUSE_FINALIZE) + | OMP_CLAUSE_COPYOUT | OMP_CLAUSE_DELETE | OMP_CLAUSE_FINALIZE \ + | OMP_CLAUSE_DETACH) #define OACC_WAIT_CLAUSES \ omp_mask (OMP_CLAUSE_ASYNC) #define OACC_ROUTINE_CLAUSES \ @@ -3812,9 +3839,6 @@ resolve_nonnegative_int_expr (gfc_expr *expr, const char *clause) static void check_symbol_not_pointer (gfc_symbol *sym, locus loc, const char *name) { - if (sym->ts.type == BT_DERIVED && sym->attr.pointer) - gfc_error ("POINTER object %qs of derived type in %s clause at %L", - sym->name, name, &loc); if (sym->ts.type == BT_DERIVED && sym->attr.cray_pointer) gfc_error ("Cray pointer object %qs of derived type in %s clause at %L", sym->name, name, &loc); @@ -3859,9 +3883,6 @@ check_array_not_assumed (gfc_symbol *sym, locus loc, const char *name) static void resolve_oacc_data_clauses (gfc_symbol *sym, locus loc, const char *name) { - if (sym->ts.type == BT_DERIVED && sym->attr.allocatable) - gfc_error ("ALLOCATABLE object %qs of derived type in %s clause at %L", - sym->name, name, &loc); if ((sym->ts.type == BT_ASSUMED && sym->attr.allocatable) || (sym->ts.type == BT_CLASS && CLASS_DATA (sym) && CLASS_DATA (sym)->attr.allocatable)) @@ -4244,11 +4265,26 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses, && (list != OMP_LIST_REDUCTION || !openacc)) for (n = omp_clauses->lists[list]; n; n = n->next) { - if (n->sym->mark) - gfc_error ("Symbol %qs present on multiple clauses at %L", - n->sym->name, &n->where); - else - n->sym->mark = 1; + bool array_only_p = true; + /* Disallow duplicate bare variable references and multiple + subarrays of the same array here, but allow multiple components of + the same (e.g. derived-type) variable. For the latter, duplicate + components are detected elsewhere. */ + if (openacc && n->expr && n->expr->expr_type == EXPR_VARIABLE) + for (gfc_ref *ref = n->expr->ref; ref; ref = ref->next) + if (ref->type != REF_ARRAY) + { + array_only_p = false; + break; + } + if (array_only_p) + { + if (n->sym->mark) + gfc_error ("Symbol %qs present on multiple clauses at %L", + n->sym->name, &n->where); + else + n->sym->mark = 1; + } } gcc_assert (OMP_LIST_LASTPRIVATE == OMP_LIST_FIRSTPRIVATE + 1); @@ -4439,23 +4475,43 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses, "are allowed on ORDERED directive at %L", &n->where); } + gfc_ref *array_ref = NULL; + bool resolved = false; if (n->expr) { - if (!gfc_resolve_expr (n->expr) + array_ref = n->expr->ref; + resolved = gfc_resolve_expr (n->expr); + + /* Look through component refs to find last array + reference. */ + if (openacc) + while (resolved + && array_ref + && (array_ref->type == REF_COMPONENT + || (array_ref->type == REF_ARRAY + && array_ref->next + && (array_ref->next->type + == REF_COMPONENT)))) + array_ref = array_ref->next; + } + if (array_ref + || (n->expr + && (!resolved || n->expr->expr_type != EXPR_VARIABLE))) + { + if (!resolved || n->expr->expr_type != EXPR_VARIABLE - || n->expr->ref == NULL - || n->expr->ref->next - || n->expr->ref->type != REF_ARRAY) + || array_ref->next + || array_ref->type != REF_ARRAY) gfc_error ("%qs in %s clause at %L is not a proper " "array section", n->sym->name, name, &n->where); - else if (n->expr->ref->u.ar.codimen) + else if (array_ref->u.ar.codimen) gfc_error ("Coarrays not supported in %s clause at %L", name, &n->where); else { int i; - gfc_array_ref *ar = &n->expr->ref->u.ar; + gfc_array_ref *ar = &array_ref->u.ar; for (i = 0; i < ar->dimen; i++) if (ar->stride[i]) { diff --git a/gcc/fortran/trans-expr.c b/gcc/fortran/trans-expr.c index 61db4e39210..b0d99f52b2e 100644 --- a/gcc/fortran/trans-expr.c +++ b/gcc/fortran/trans-expr.c @@ -2422,7 +2422,7 @@ gfc_conv_substring (gfc_se * se, gfc_ref * ref, int kind, /* Convert a derived type component reference. */ -static void +void gfc_conv_component_ref (gfc_se * se, gfc_ref * ref) { gfc_component *c; @@ -2512,7 +2512,7 @@ gfc_conv_component_ref (gfc_se * se, gfc_ref * ref) /* This function deals with component references to components of the parent type for derived type extensions. */ -static void +void conv_parent_component_references (gfc_se * se, gfc_ref * ref) { gfc_component *c; @@ -2578,6 +2578,95 @@ conv_inquiry (gfc_se * se, gfc_ref * ref, gfc_expr *expr, gfc_typespec *ts) se->expr = res; } +/* Transparently dereference VAR if it is a pointer, reference, etc. + according to Fortran semantics. */ + +tree +gfc_auto_dereference_var (location_t loc, gfc_symbol *sym, tree var, + bool descriptor_only_p, bool is_classarray) +{ + /* Characters are entirely different from other types, they are treated + separately. */ + if (sym->ts.type == BT_CHARACTER) + { + /* Dereference character pointer dummy arguments + or results. */ + if ((sym->attr.pointer || sym->attr.allocatable) + && (sym->attr.dummy + || sym->attr.function + || sym->attr.result)) + var = build_fold_indirect_ref_loc (input_location, var); + } + else if (!sym->attr.value) + { + /* Dereference temporaries for class array dummy arguments. */ + if (sym->attr.dummy && is_classarray + && GFC_ARRAY_TYPE_P (TREE_TYPE (var))) + { + if (!descriptor_only_p) + var = GFC_DECL_SAVED_DESCRIPTOR (var); + + var = build_fold_indirect_ref_loc (input_location, var); + } + + /* Dereference non-character scalar dummy arguments. */ + if (sym->attr.dummy && !sym->attr.dimension + && !(sym->attr.codimension && sym->attr.allocatable) + && (sym->ts.type != BT_CLASS + || (!CLASS_DATA (sym)->attr.dimension + && !(CLASS_DATA (sym)->attr.codimension + && CLASS_DATA (sym)->attr.allocatable)))) + var = build_fold_indirect_ref_loc (input_location, var); + + /* Dereference scalar hidden result. */ + if (flag_f2c && sym->ts.type == BT_COMPLEX + && (sym->attr.function || sym->attr.result) + && !sym->attr.dimension && !sym->attr.pointer + && !sym->attr.always_explicit) + var = build_fold_indirect_ref_loc (input_location, var); + + /* Dereference non-character, non-class pointer variables. + These must be dummies, results, or scalars. */ + if (!is_classarray + && (sym->attr.pointer || sym->attr.allocatable + || gfc_is_associate_pointer (sym) + || (sym->as && sym->as->type == AS_ASSUMED_RANK)) + && (sym->attr.dummy + || sym->attr.function + || sym->attr.result + || (!sym->attr.dimension + && (!sym->attr.codimension || !sym->attr.allocatable)))) + var = build_fold_indirect_ref_loc (input_location, var); + /* Now treat the class array pointer variables accordingly. */ + else if (sym->ts.type == BT_CLASS + && sym->attr.dummy + && (CLASS_DATA (sym)->attr.dimension + || CLASS_DATA (sym)->attr.codimension) + && ((CLASS_DATA (sym)->as + && CLASS_DATA (sym)->as->type == AS_ASSUMED_RANK) + || CLASS_DATA (sym)->attr.allocatable + || CLASS_DATA (sym)->attr.class_pointer)) + var = build_fold_indirect_ref_loc (input_location, var); + /* And the case where a non-dummy, non-result, non-function, + non-allotable and non-pointer classarray is present. This case was + previously covered by the first if, but with introducing the + condition !is_classarray there, that case has to be covered + explicitly. */ + else if (sym->ts.type == BT_CLASS + && !sym->attr.dummy + && !sym->attr.function + && !sym->attr.result + && (CLASS_DATA (sym)->attr.dimension + || CLASS_DATA (sym)->attr.codimension) + && (sym->assoc + || !CLASS_DATA (sym)->attr.allocatable) + && !CLASS_DATA (sym)->attr.class_pointer) + var = build_fold_indirect_ref_loc (input_location, var); + } + + return var; +} + /* Return the contents of a variable. Also handles reference/pointer variables (all Fortran pointer references are implicit). */ @@ -2684,94 +2773,9 @@ gfc_conv_variable (gfc_se * se, gfc_expr * expr) return; } - - /* Dereference the expression, where needed. Since characters - are entirely different from other types, they are treated - separately. */ - if (sym->ts.type == BT_CHARACTER) - { - /* Dereference character pointer dummy arguments - or results. */ - if ((sym->attr.pointer || sym->attr.allocatable) - && (sym->attr.dummy - || sym->attr.function - || sym->attr.result)) - se->expr = build_fold_indirect_ref_loc (input_location, - se->expr); - - } - else if (!sym->attr.value) - { - /* Dereference temporaries for class array dummy arguments. */ - if (sym->attr.dummy && is_classarray - && GFC_ARRAY_TYPE_P (TREE_TYPE (se->expr))) - { - if (!se->descriptor_only) - se->expr = GFC_DECL_SAVED_DESCRIPTOR (se->expr); - - se->expr = build_fold_indirect_ref_loc (input_location, - se->expr); - } - - /* Dereference non-character scalar dummy arguments. */ - if (sym->attr.dummy && !sym->attr.dimension - && !(sym->attr.codimension && sym->attr.allocatable) - && (sym->ts.type != BT_CLASS - || (!CLASS_DATA (sym)->attr.dimension - && !(CLASS_DATA (sym)->attr.codimension - && CLASS_DATA (sym)->attr.allocatable)))) - se->expr = build_fold_indirect_ref_loc (input_location, - se->expr); - - /* Dereference scalar hidden result. */ - if (flag_f2c && sym->ts.type == BT_COMPLEX - && (sym->attr.function || sym->attr.result) - && !sym->attr.dimension && !sym->attr.pointer - && !sym->attr.always_explicit) - se->expr = build_fold_indirect_ref_loc (input_location, - se->expr); - - /* Dereference non-character, non-class pointer variables. - These must be dummies, results, or scalars. */ - if (!is_classarray - && (sym->attr.pointer || sym->attr.allocatable - || gfc_is_associate_pointer (sym) - || (sym->as && sym->as->type == AS_ASSUMED_RANK)) - && (sym->attr.dummy - || sym->attr.function - || sym->attr.result - || (!sym->attr.dimension - && (!sym->attr.codimension || !sym->attr.allocatable)))) - se->expr = build_fold_indirect_ref_loc (input_location, - se->expr); - /* Now treat the class array pointer variables accordingly. */ - else if (sym->ts.type == BT_CLASS - && sym->attr.dummy - && (CLASS_DATA (sym)->attr.dimension - || CLASS_DATA (sym)->attr.codimension) - && ((CLASS_DATA (sym)->as - && CLASS_DATA (sym)->as->type == AS_ASSUMED_RANK) - || CLASS_DATA (sym)->attr.allocatable - || CLASS_DATA (sym)->attr.class_pointer)) - se->expr = build_fold_indirect_ref_loc (input_location, - se->expr); - /* And the case where a non-dummy, non-result, non-function, - non-allotable and non-pointer classarray is present. This case was - previously covered by the first if, but with introducing the - condition !is_classarray there, that case has to be covered - explicitly. */ - else if (sym->ts.type == BT_CLASS - && !sym->attr.dummy - && !sym->attr.function - && !sym->attr.result - && (CLASS_DATA (sym)->attr.dimension - || CLASS_DATA (sym)->attr.codimension) - && (sym->assoc - || !CLASS_DATA (sym)->attr.allocatable) - && !CLASS_DATA (sym)->attr.class_pointer) - se->expr = build_fold_indirect_ref_loc (input_location, - se->expr); - } + /* Dereference the expression, where needed. */ + se->expr = gfc_auto_dereference_var (input_location, sym, se->expr, + se->descriptor_only, is_classarray); ref = expr->ref; } diff --git a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c index f83bab4850e..775548fe9af 100644 --- a/gcc/fortran/trans-openmp.c +++ b/gcc/fortran/trans-openmp.c @@ -75,6 +75,9 @@ gfc_omp_privatize_by_reference (const_tree decl) if (TREE_CODE (type) == POINTER_TYPE) { + while (TREE_CODE (decl) == COMPONENT_REF) + decl = TREE_OPERAND (decl, 1); + /* Array POINTER/ALLOCATABLE have aggregate types, all user variables that have POINTER_TYPE type and aren't scalar pointers, scalar allocatables, Cray pointees or C pointers are supposed to be @@ -1836,6 +1839,92 @@ gfc_convert_expr_to_tree (stmtblock_t *block, gfc_expr *expr) static vec *doacross_steps; + +/* Translate an array section or array element. */ + +static void +gfc_trans_omp_array_section (stmtblock_t *block, gfc_omp_namelist *n, + tree decl, bool element, gomp_map_kind ptr_kind, + tree node, tree &node2, tree &node3, tree &node4) +{ + gfc_se se; + tree ptr, ptr2; + + gfc_init_se (&se, NULL); + + if (element) + { + gfc_conv_expr_reference (&se, n->expr); + gfc_add_block_to_block (block, &se.pre); + ptr = se.expr; + OMP_CLAUSE_SIZE (node) + = TYPE_SIZE_UNIT (TREE_TYPE (ptr)); + } + else + { + gfc_conv_expr_descriptor (&se, n->expr); + ptr = gfc_conv_array_data (se.expr); + tree type = TREE_TYPE (se.expr); + gfc_add_block_to_block (block, &se.pre); + OMP_CLAUSE_SIZE (node) = gfc_full_array_size (block, se.expr, + GFC_TYPE_ARRAY_RANK (type)); + tree elemsz = TYPE_SIZE_UNIT (gfc_get_element_type (type)); + elemsz = fold_convert (gfc_array_index_type, elemsz); + OMP_CLAUSE_SIZE (node) = fold_build2 (MULT_EXPR, gfc_array_index_type, + OMP_CLAUSE_SIZE (node), elemsz); + } + gfc_add_block_to_block (block, &se.post); + ptr = fold_convert (build_pointer_type (char_type_node), ptr); + OMP_CLAUSE_DECL (node) = build_fold_indirect_ref (ptr); + + if (POINTER_TYPE_P (TREE_TYPE (decl)) + && GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (TREE_TYPE (decl))) + && ptr_kind == GOMP_MAP_POINTER) + { + node4 = build_omp_clause (input_location, + OMP_CLAUSE_MAP); + OMP_CLAUSE_SET_MAP_KIND (node4, GOMP_MAP_POINTER); + OMP_CLAUSE_DECL (node4) = decl; + OMP_CLAUSE_SIZE (node4) = size_int (0); + decl = build_fold_indirect_ref (decl); + } + ptr = fold_convert (sizetype, ptr); + if (GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (decl))) + { + tree type = TREE_TYPE (decl); + ptr2 = gfc_conv_descriptor_data_get (decl); + node2 = build_omp_clause (input_location, + OMP_CLAUSE_MAP); + OMP_CLAUSE_SET_MAP_KIND (node2, GOMP_MAP_TO_PSET); + OMP_CLAUSE_DECL (node2) = decl; + OMP_CLAUSE_SIZE (node2) = TYPE_SIZE_UNIT (type); + node3 = build_omp_clause (input_location, + OMP_CLAUSE_MAP); + OMP_CLAUSE_SET_MAP_KIND (node3, ptr_kind); + OMP_CLAUSE_DECL (node3) + = gfc_conv_descriptor_data_get (decl); + if (ptr_kind == GOMP_MAP_ATTACH_DETACH) + STRIP_NOPS (OMP_CLAUSE_DECL (node3)); + } + else + { + if (TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE) + ptr2 = build_fold_addr_expr (decl); + else + { + gcc_assert (POINTER_TYPE_P (TREE_TYPE (decl))); + ptr2 = decl; + } + node3 = build_omp_clause (input_location, + OMP_CLAUSE_MAP); + OMP_CLAUSE_SET_MAP_KIND (node3, ptr_kind); + OMP_CLAUSE_DECL (node3) = decl; + } + ptr2 = fold_convert (sizetype, ptr2); + OMP_CLAUSE_SIZE (node3) + = fold_build2 (MINUS_EXPR, sizetype, ptr, ptr2); +} + static tree gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, locus where, bool declare_simd = false) @@ -2161,7 +2250,8 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, || GFC_DECL_GET_SCALAR_ALLOCATABLE (decl) || GFC_DECL_CRAY_POINTEE (decl) || GFC_DESCRIPTOR_TYPE_P - (TREE_TYPE (TREE_TYPE (decl))))) + (TREE_TYPE (TREE_TYPE (decl))) + || n->sym->ts.type == BT_DERIVED)) { tree orig_decl = decl; node4 = build_omp_clause (input_location, @@ -2182,10 +2272,13 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, decl = build_fold_indirect_ref (decl); } } - if (GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (decl))) + if (GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (decl)) + && n->u.map_op != OMP_MAP_ATTACH + && n->u.map_op != OMP_MAP_DETACH) { tree type = TREE_TYPE (decl); tree ptr = gfc_conv_descriptor_data_get (decl); + ptr = fold_convert (build_pointer_type (char_type_node), ptr); ptr = build_fold_indirect_ref (ptr); @@ -2251,88 +2344,154 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, else OMP_CLAUSE_DECL (node) = decl; } - else + else if (n->expr + && n->expr->expr_type == EXPR_VARIABLE + && n->expr->ref->type == REF_COMPONENT) { - tree ptr, ptr2; + gfc_ref *lastcomp; + + for (gfc_ref *ref = n->expr->ref; ref; ref = ref->next) + if (ref->type == REF_COMPONENT) + lastcomp = ref; + + symbol_attribute sym_attr; + + sym_attr = lastcomp->u.c.component->attr; + gfc_init_se (&se, NULL); - if (n->expr->ref->u.ar.type == AR_ELEMENT) - { - gfc_conv_expr_reference (&se, n->expr); - gfc_add_block_to_block (block, &se.pre); - ptr = se.expr; - OMP_CLAUSE_SIZE (node) - = TYPE_SIZE_UNIT (TREE_TYPE (ptr)); - } - else + + if (!sym_attr.dimension + && lastcomp->u.c.component->ts.type != BT_DERIVED) { - gfc_conv_expr_descriptor (&se, n->expr); - ptr = gfc_conv_array_data (se.expr); - tree type = TREE_TYPE (se.expr); + /* Last component is a scalar. */ + gfc_conv_expr (&se, n->expr); gfc_add_block_to_block (block, &se.pre); - OMP_CLAUSE_SIZE (node) - = gfc_full_array_size (block, se.expr, - GFC_TYPE_ARRAY_RANK (type)); - tree elemsz - = TYPE_SIZE_UNIT (gfc_get_element_type (type)); - elemsz = fold_convert (gfc_array_index_type, elemsz); - OMP_CLAUSE_SIZE (node) - = fold_build2 (MULT_EXPR, gfc_array_index_type, - OMP_CLAUSE_SIZE (node), elemsz); + OMP_CLAUSE_DECL (node) = se.expr; + gfc_add_block_to_block (block, &se.post); + goto finalize_map_clause; } - gfc_add_block_to_block (block, &se.post); - ptr = fold_convert (build_pointer_type (char_type_node), - ptr); - OMP_CLAUSE_DECL (node) = build_fold_indirect_ref (ptr); - if (POINTER_TYPE_P (TREE_TYPE (decl)) - && GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (TREE_TYPE (decl)))) + se.expr + = gfc_auto_dereference_var (input_location, n->sym, + decl); + + for (gfc_ref *ref = n->expr->ref; + ref && ref != lastcomp->next; + ref = ref->next) { - node4 = build_omp_clause (input_location, - OMP_CLAUSE_MAP); - OMP_CLAUSE_SET_MAP_KIND (node4, GOMP_MAP_POINTER); - OMP_CLAUSE_DECL (node4) = decl; - OMP_CLAUSE_SIZE (node4) = size_int (0); - decl = build_fold_indirect_ref (decl); + if (ref->type == REF_COMPONENT) + { + if (ref->u.c.sym->attr.extension) + conv_parent_component_references (&se, ref); + + gfc_conv_component_ref (&se, ref); + } + else + sorry ("unhandled derived-type component"); } - ptr = fold_convert (sizetype, ptr); - if (GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (decl))) + + tree inner = se.expr; + + /* Last component is a derived type. */ + if (lastcomp->u.c.component->ts.type == BT_DERIVED) { - tree type = TREE_TYPE (decl); - ptr2 = gfc_conv_descriptor_data_get (decl); - node2 = build_omp_clause (input_location, - OMP_CLAUSE_MAP); - OMP_CLAUSE_SET_MAP_KIND (node2, GOMP_MAP_TO_PSET); - OMP_CLAUSE_DECL (node2) = decl; - OMP_CLAUSE_SIZE (node2) = TYPE_SIZE_UNIT (type); - node3 = build_omp_clause (input_location, - OMP_CLAUSE_MAP); - OMP_CLAUSE_SET_MAP_KIND (node3, GOMP_MAP_POINTER); - OMP_CLAUSE_DECL (node3) - = gfc_conv_descriptor_data_get (decl); + if (sym_attr.allocatable || sym_attr.pointer) + { + tree data = inner; + tree size = TYPE_SIZE_UNIT (TREE_TYPE (inner)); + + OMP_CLAUSE_DECL (node) + = build_fold_indirect_ref (data); + OMP_CLAUSE_SIZE (node) = size; + node2 = build_omp_clause (input_location, + OMP_CLAUSE_MAP); + OMP_CLAUSE_SET_MAP_KIND (node2, + GOMP_MAP_ATTACH_DETACH); + OMP_CLAUSE_DECL (node2) = data; + OMP_CLAUSE_SIZE (node2) = size_int (0); + } + else + { + OMP_CLAUSE_DECL (node) = decl; + OMP_CLAUSE_SIZE (node) + = TYPE_SIZE_UNIT (TREE_TYPE (decl)); + } } - else + else if (lastcomp->next + && lastcomp->next->type == REF_ARRAY + && lastcomp->next->u.ar.type == AR_FULL) { - if (TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE) - ptr2 = build_fold_addr_expr (decl); - else + /* Just pass the (auto-dereferenced) decl through for + bare attach and detach clauses. */ + if (n->u.map_op == OMP_MAP_ATTACH + || n->u.map_op == OMP_MAP_DETACH) { - gcc_assert (POINTER_TYPE_P (TREE_TYPE (decl))); - ptr2 = decl; + OMP_CLAUSE_DECL (node) = inner; + OMP_CLAUSE_SIZE (node) = size_zero_node; + goto finalize_map_clause; } - node3 = build_omp_clause (input_location, - OMP_CLAUSE_MAP); - OMP_CLAUSE_SET_MAP_KIND (node3, GOMP_MAP_POINTER); - OMP_CLAUSE_DECL (node3) = decl; + + if (GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (inner))) + { + tree type = TREE_TYPE (inner); + tree ptr = gfc_conv_descriptor_data_get (inner); + ptr = build_fold_indirect_ref (ptr); + OMP_CLAUSE_DECL (node) = ptr; + node2 = build_omp_clause (input_location, + OMP_CLAUSE_MAP); + OMP_CLAUSE_SET_MAP_KIND (node2, GOMP_MAP_TO_PSET); + OMP_CLAUSE_DECL (node2) = inner; + OMP_CLAUSE_SIZE (node2) = TYPE_SIZE_UNIT (type); + node3 = build_omp_clause (input_location, + OMP_CLAUSE_MAP); + OMP_CLAUSE_SET_MAP_KIND (node3, + GOMP_MAP_ATTACH_DETACH); + OMP_CLAUSE_DECL (node3) + = gfc_conv_descriptor_data_get (inner); + STRIP_NOPS (OMP_CLAUSE_DECL (node3)); + OMP_CLAUSE_SIZE (node3) = size_int (0); + int rank = GFC_TYPE_ARRAY_RANK (type); + OMP_CLAUSE_SIZE (node) + = gfc_full_array_size (block, inner, rank); + tree elemsz + = TYPE_SIZE_UNIT (gfc_get_element_type (type)); + elemsz = fold_convert (gfc_array_index_type, elemsz); + OMP_CLAUSE_SIZE (node) + = fold_build2 (MULT_EXPR, gfc_array_index_type, + OMP_CLAUSE_SIZE (node), elemsz); + } + else + OMP_CLAUSE_DECL (node) = inner; + } + else /* An array element or section. */ + { + bool element + = (lastcomp->next + && lastcomp->next->type == REF_ARRAY + && lastcomp->next->u.ar.type == AR_ELEMENT); + + gfc_trans_omp_array_section (block, n, inner, element, + GOMP_MAP_ATTACH_DETACH, + node, node2, node3, node4); } - ptr2 = fold_convert (sizetype, ptr2); - OMP_CLAUSE_SIZE (node3) - = fold_build2 (MINUS_EXPR, sizetype, ptr, ptr2); } + else /* An array element or array section. */ + { + bool element = n->expr->ref->u.ar.type == AR_ELEMENT; + gfc_trans_omp_array_section (block, n, decl, element, + GOMP_MAP_POINTER, node, node2, + node3, node4); + } + + finalize_map_clause: switch (n->u.map_op) { case OMP_MAP_ALLOC: OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_ALLOC); break; + case OMP_MAP_ATTACH: + OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_ATTACH); + break; case OMP_MAP_TO: OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_TO); break; @@ -2357,6 +2516,9 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, case OMP_MAP_DELETE: OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_DELETE); break; + case OMP_MAP_DETACH: + OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_DETACH); + break; case OMP_MAP_FORCE_ALLOC: OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_FORCE_ALLOC); break; diff --git a/gcc/fortran/trans.h b/gcc/fortran/trans.h index 405e88dd1c4..b9999d0d2ce 100644 --- a/gcc/fortran/trans.h +++ b/gcc/fortran/trans.h @@ -565,6 +565,15 @@ tree gfc_conv_expr_present (gfc_symbol *); /* Convert a missing, dummy argument into a null or zero. */ void gfc_conv_missing_dummy (gfc_se *, gfc_expr *, gfc_typespec, int); +/* Lowering of component references. */ +void gfc_conv_component_ref (gfc_se * se, gfc_ref * ref); +void conv_parent_component_references (gfc_se * se, gfc_ref * ref); + +/* Automatically dereference var. */ +tree gfc_auto_dereference_var (location_t, gfc_symbol *, tree, + bool desc_only = false, + bool is_classarray = false); + /* Generate code to allocate a string temporary. */ tree gfc_conv_string_tmp (gfc_se *, tree, tree); /* Get the string length variable belonging to an expression. */ diff --git a/gcc/gimplify.c b/gcc/gimplify.c index 5d50713a9f1..421c6cb7894 100644 --- a/gcc/gimplify.c +++ b/gcc/gimplify.c @@ -123,6 +123,10 @@ enum gimplify_omp_var_data /* Flag for GOVD_REDUCTION: inscan seen in {in,ex}clusive clause. */ GOVD_REDUCTION_INSCAN = 0x2000000, + /* Flag for GOVD_MAP: (struct) vars that have pointer attachments for + fields. */ + GOVD_MAP_HAS_ATTACHMENTS = 8388608, + GOVD_DATA_SHARE_CLASS = (GOVD_SHARED | GOVD_PRIVATE | GOVD_FIRSTPRIVATE | GOVD_LASTPRIVATE | GOVD_REDUCTION | GOVD_LINEAR | GOVD_LOCAL) @@ -8151,7 +8155,13 @@ insert_struct_comp_map (enum tree_code code, tree c, tree struct_node, OMP_CLAUSE_SET_MAP_KIND (c2, mkind); OMP_CLAUSE_DECL (c2) = unshare_expr (OMP_CLAUSE_DECL (c)); OMP_CLAUSE_CHAIN (c2) = scp ? *scp : prev_node; - OMP_CLAUSE_SIZE (c2) = TYPE_SIZE_UNIT (ptr_type_node); + if (OMP_CLAUSE_CHAIN (prev_node) != c + && OMP_CLAUSE_CODE (OMP_CLAUSE_CHAIN (prev_node)) == OMP_CLAUSE_MAP + && (OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (prev_node)) + == GOMP_MAP_TO_PSET)) + OMP_CLAUSE_SIZE (c2) = OMP_CLAUSE_SIZE (OMP_CLAUSE_CHAIN (prev_node)); + else + OMP_CLAUSE_SIZE (c2) = TYPE_SIZE_UNIT (ptr_type_node); if (struct_node) OMP_CLAUSE_CHAIN (struct_node) = c2; @@ -8160,8 +8170,10 @@ insert_struct_comp_map (enum tree_code code, tree c, tree struct_node, GOMP_MAP_ALWAYS_POINTER though, i.e. a GOMP_MAP_TO_PSET. */ if (OMP_CLAUSE_CHAIN (prev_node) != c && OMP_CLAUSE_CODE (OMP_CLAUSE_CHAIN (prev_node)) == OMP_CLAUSE_MAP - && (OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (prev_node)) - == GOMP_MAP_ALWAYS_POINTER)) + && ((OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (prev_node)) + == GOMP_MAP_ALWAYS_POINTER) + || (OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (prev_node)) + == GOMP_MAP_ATTACH_DETACH))) { tree c4 = OMP_CLAUSE_CHAIN (prev_node); tree c3 = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP); @@ -8285,6 +8297,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, struct gimplify_omp_ctx *ctx, *outer_ctx; tree c; hash_map *struct_map_to_clause = NULL; + hash_set *struct_deref_set = NULL; tree *prev_list_p = NULL, *orig_list_p = list_p; int handled_depend_iterators = -1; int nowait = -1; @@ -8686,8 +8699,6 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, case OMP_TARGET_DATA: case OMP_TARGET_ENTER_DATA: case OMP_TARGET_EXIT_DATA: - case OACC_ENTER_DATA: - case OACC_EXIT_DATA: case OACC_HOST_DATA: if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER || (OMP_CLAUSE_MAP_KIND (c) @@ -8696,6 +8707,15 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, mapped, but not the pointer to it. */ remove = true; break; + case OACC_ENTER_DATA: + case OACC_EXIT_DATA: + if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER + || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_TO_PSET + || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER + || (OMP_CLAUSE_MAP_KIND (c) + == GOMP_MAP_FIRSTPRIVATE_REFERENCE)) + remove = true; + break; default: break; } @@ -8758,7 +8778,35 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, pd = &TREE_OPERAND (decl, 0); decl = TREE_OPERAND (decl, 0); } - if (TREE_CODE (decl) == COMPONENT_REF) + bool indir_p = false; + tree orig_decl = decl; + tree decl_ref = NULL_TREE; + if ((region_type & ORT_ACC) != 0 + && TREE_CODE (*pd) == COMPONENT_REF + && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH + && code != OACC_UPDATE) + { + while (TREE_CODE (decl) == COMPONENT_REF) + { + decl = TREE_OPERAND (decl, 0); + if ((TREE_CODE (decl) == MEM_REF + && integer_zerop (TREE_OPERAND (decl, 1))) + || INDIRECT_REF_P (decl)) + { + indir_p = true; + decl = TREE_OPERAND (decl, 0); + } + if (TREE_CODE (decl) == INDIRECT_REF + && DECL_P (TREE_OPERAND (decl, 0)) + && (TREE_CODE (TREE_TYPE (TREE_OPERAND (decl, 0))) + == REFERENCE_TYPE)) + { + decl_ref = decl; + decl = TREE_OPERAND (decl, 0); + } + } + } + else if (TREE_CODE (decl) == COMPONENT_REF) { while (TREE_CODE (decl) == COMPONENT_REF) decl = TREE_OPERAND (decl, 0); @@ -8768,13 +8816,76 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, == REFERENCE_TYPE)) decl = TREE_OPERAND (decl, 0); } + if (decl != orig_decl && DECL_P (decl) && indir_p) + { + gomp_map_kind k = (code == OACC_EXIT_DATA) ? GOMP_MAP_DETACH + : GOMP_MAP_ATTACH; + /* We have a dereference of a struct member. Make this an + attach/detach operation, and ensure the base pointer is + mapped as a FIRSTPRIVATE_POINTER. */ + OMP_CLAUSE_SET_MAP_KIND (c, k); + flags = GOVD_MAP | GOVD_SEEN | GOVD_EXPLICIT; + tree next_clause = OMP_CLAUSE_CHAIN (c); + if (k == GOMP_MAP_ATTACH + && code != OACC_ENTER_DATA + && (!next_clause + || (OMP_CLAUSE_CODE (next_clause) != OMP_CLAUSE_MAP) + || (OMP_CLAUSE_MAP_KIND (next_clause) + != GOMP_MAP_POINTER) + || OMP_CLAUSE_DECL (next_clause) != decl) + && (!struct_deref_set + || !struct_deref_set->contains (decl))) + { + if (!struct_deref_set) + struct_deref_set = new hash_set (); + /* As well as the attach, we also need a + FIRSTPRIVATE_POINTER clause to properly map the + pointer to the struct base. */ + tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c), + OMP_CLAUSE_MAP); + OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ALLOC); + OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION (c2) + = 1; + tree charptr_zero + = build_int_cst (build_pointer_type (char_type_node), + 0); + OMP_CLAUSE_DECL (c2) + = build2 (MEM_REF, char_type_node, + decl_ref ? decl_ref : decl, charptr_zero); + OMP_CLAUSE_SIZE (c2) = size_zero_node; + tree c3 = build_omp_clause (OMP_CLAUSE_LOCATION (c), + OMP_CLAUSE_MAP); + OMP_CLAUSE_SET_MAP_KIND (c3, + GOMP_MAP_FIRSTPRIVATE_POINTER); + OMP_CLAUSE_DECL (c3) = decl; + OMP_CLAUSE_SIZE (c3) = size_zero_node; + tree mapgrp = *prev_list_p; + *prev_list_p = c2; + OMP_CLAUSE_CHAIN (c3) = mapgrp; + OMP_CLAUSE_CHAIN (c2) = c3; + + struct_deref_set->add (decl); + } + goto do_add_decl; + } + /* An "attach/detach" operation on an update directive should + behave as a GOMP_MAP_ALWAYS_POINTER. Beware that + unlike attach or detach map kinds, GOMP_MAP_ALWAYS_POINTER + depends on the previous mapping. */ + if (code == OACC_UPDATE + && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH) + OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_ALWAYS_POINTER); if (gimplify_expr (pd, pre_p, NULL, is_gimple_lvalue, fb_lvalue) == GS_ERROR) { remove = true; break; } - if (DECL_P (decl)) + if (DECL_P (decl) + && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_TO_PSET + && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ATTACH + && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_DETACH + && code != OACC_UPDATE) { if (error_operand_p (decl)) { @@ -8795,7 +8906,8 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, break; } - if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_POINTER) + if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_POINTER + || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH) { /* Error recovery. */ if (prev_list_p == NULL) @@ -8827,20 +8939,46 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, = splay_tree_lookup (ctx->variables, (splay_tree_key)decl); bool ptr = (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_POINTER); + bool attach_detach = (OMP_CLAUSE_MAP_KIND (c) + == GOMP_MAP_ATTACH_DETACH); + bool attach = OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH + || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH; + bool has_attachments = false; + /* For OpenACC, pointers in structs should trigger an + attach action. */ + if (attach_detach && (region_type & ORT_ACC) != 0) + { + /* Turning a GOMP_MAP_ALWAYS_POINTER clause into a + GOMP_MAP_ATTACH clause after we have detected a case + that needs a GOMP_MAP_STRUCT mapping added. */ + gomp_map_kind k + = (code == OACC_EXIT_DATA) ? GOMP_MAP_DETACH + : GOMP_MAP_ATTACH; + OMP_CLAUSE_SET_MAP_KIND (c, k); + has_attachments = true; + } if (n == NULL || (n->value & GOVD_MAP) == 0) { tree l = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP); - OMP_CLAUSE_SET_MAP_KIND (l, GOMP_MAP_STRUCT); + gomp_map_kind k = attach ? GOMP_MAP_FORCE_PRESENT + : GOMP_MAP_STRUCT; + + OMP_CLAUSE_SET_MAP_KIND (l, k); if (!base_eq_orig_base) OMP_CLAUSE_DECL (l) = unshare_expr (orig_base); else OMP_CLAUSE_DECL (l) = decl; - OMP_CLAUSE_SIZE (l) = size_int (1); + OMP_CLAUSE_SIZE (l) + = (!attach + ? size_int (1) + : DECL_P (OMP_CLAUSE_DECL (l)) + ? DECL_SIZE_UNIT (OMP_CLAUSE_DECL (l)) + : TYPE_SIZE_UNIT (TREE_TYPE (OMP_CLAUSE_DECL (l)))); if (struct_map_to_clause == NULL) struct_map_to_clause = new hash_map; struct_map_to_clause->put (decl, l); - if (ptr) + if (ptr || attach_detach) { insert_struct_comp_map (code, c, l, *prev_list_p, NULL); @@ -8866,23 +9004,31 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, OMP_CLAUSE_CHAIN (l) = c2; } flags = GOVD_MAP | GOVD_EXPLICIT; - if (GOMP_MAP_ALWAYS_P (OMP_CLAUSE_MAP_KIND (c)) || ptr) + if (GOMP_MAP_ALWAYS_P (OMP_CLAUSE_MAP_KIND (c)) + || ptr + || attach_detach) flags |= GOVD_SEEN; + if (has_attachments) + flags |= GOVD_MAP_HAS_ATTACHMENTS; goto do_add_decl; } - else + else if (struct_map_to_clause) { tree *osc = struct_map_to_clause->get (decl); tree *sc = NULL, *scp = NULL; - if (GOMP_MAP_ALWAYS_P (OMP_CLAUSE_MAP_KIND (c)) || ptr) + if (GOMP_MAP_ALWAYS_P (OMP_CLAUSE_MAP_KIND (c)) + || ptr + || attach_detach) n->value |= GOVD_SEEN; sc = &OMP_CLAUSE_CHAIN (*osc); if (*sc != c && (OMP_CLAUSE_MAP_KIND (*sc) - == GOMP_MAP_FIRSTPRIVATE_REFERENCE)) + == GOMP_MAP_FIRSTPRIVATE_REFERENCE)) sc = &OMP_CLAUSE_CHAIN (*sc); + /* Here "prev_list_p" is the end of the inserted + alloc/release nodes after the struct node, OSC. */ for (; *sc != c; sc = &OMP_CLAUSE_CHAIN (*sc)) - if (ptr && sc == prev_list_p) + if ((ptr || attach_detach) && sc == prev_list_p) break; else if (TREE_CODE (OMP_CLAUSE_DECL (*sc)) != COMPONENT_REF @@ -8931,7 +9077,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, } if (same_decl_offset_lt) { - if (ptr) + if (ptr || attach_detach) scp = sc; else break; @@ -8939,10 +9085,11 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, } if (remove) break; - OMP_CLAUSE_SIZE (*osc) - = size_binop (PLUS_EXPR, OMP_CLAUSE_SIZE (*osc), - size_one_node); - if (ptr) + if (!attach) + OMP_CLAUSE_SIZE (*osc) + = size_binop (PLUS_EXPR, OMP_CLAUSE_SIZE (*osc), + size_one_node); + if (ptr || attach_detach) { tree cl = insert_struct_comp_map (code, c, NULL, *prev_list_p, scp); @@ -8972,11 +9119,18 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, } if (!remove && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ALWAYS_POINTER + && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ATTACH_DETACH + && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_TO_PSET && OMP_CLAUSE_CHAIN (c) && OMP_CLAUSE_CODE (OMP_CLAUSE_CHAIN (c)) == OMP_CLAUSE_MAP - && (OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (c)) - == GOMP_MAP_ALWAYS_POINTER)) + && ((OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (c)) + == GOMP_MAP_ALWAYS_POINTER) + || (OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (c)) + == GOMP_MAP_ATTACH_DETACH) + || (OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (c)) + == GOMP_MAP_TO_PSET))) prev_list_p = list_p; + break; } flags = GOVD_MAP | GOVD_EXPLICIT; @@ -9500,6 +9654,8 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, gimplify_omp_ctxp = ctx; if (struct_map_to_clause) delete struct_map_to_clause; + if (struct_deref_set) + delete struct_deref_set; } /* Return true if DECL is a candidate for shared to firstprivate @@ -9647,6 +9803,8 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n, void *data) return 0; if ((flags & GOVD_SEEN) == 0) return 0; + if ((flags & GOVD_MAP_HAS_ATTACHMENTS) != 0) + return 0; if (flags & GOVD_DEBUG_PRIVATE) { gcc_assert ((flags & GOVD_DATA_SHARE_CLASS) == GOVD_SHARED); @@ -12550,8 +12708,9 @@ gimplify_omp_target_update (tree *expr_p, gimple_seq *pre_p) && omp_find_clause (OMP_STANDALONE_CLAUSES (expr), OMP_CLAUSE_FINALIZE)) { - /* Use GOMP_MAP_DELETE/GOMP_MAP_FORCE_FROM to denote that "finalize" - semantics apply to all mappings of this OpenACC directive. */ + /* Use GOMP_MAP_DELETE, GOMP_MAP_FORCE_DETACH, and + GOMP_MAP_FORCE_FROM to denote that "finalize" semantics apply + to all mappings of this OpenACC directive. */ bool finalize_marked = false; for (tree c = OMP_STANDALONE_CLAUSES (expr); c; c = OMP_CLAUSE_CHAIN (c)) if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP) @@ -12565,10 +12724,19 @@ gimplify_omp_target_update (tree *expr_p, gimple_seq *pre_p) OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_DELETE); finalize_marked = true; break; + case GOMP_MAP_DETACH: + OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_FORCE_DETACH); + finalize_marked = true; + break; + case GOMP_MAP_STRUCT: + case GOMP_MAP_FORCE_PRESENT: + /* Skip over an initial struct or force_present mapping. */ + break; default: - /* Check consistency: libgomp relies on the very first data - mapping clause being marked, so make sure we did that before - any other mapping clauses. */ + /* Check consistency: libgomp relies on the very first + non-struct, non-force-present data mapping clause being + marked, so make sure we did that before any other mapping + clauses. */ gcc_assert (finalize_marked); break; } diff --git a/gcc/omp-low.c b/gcc/omp-low.c index ca7dfdb83a1..59c6678e2dd 100644 --- a/gcc/omp-low.c +++ b/gcc/omp-low.c @@ -11322,6 +11322,9 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) case GOMP_MAP_FORCE_DEVICEPTR: case GOMP_MAP_DEVICE_RESIDENT: case GOMP_MAP_LINK: + case GOMP_MAP_ATTACH: + case GOMP_MAP_DETACH: + case GOMP_MAP_FORCE_DETACH: gcc_assert (is_gimple_omp_oacc (stmt)); break; default: diff --git a/gcc/testsuite/c-c++-common/goacc/deep-copy-arrayofstruct.c b/gcc/testsuite/c-c++-common/goacc/deep-copy-arrayofstruct.c new file mode 100644 index 00000000000..d411bcfa8e7 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/deep-copy-arrayofstruct.c @@ -0,0 +1,84 @@ +/* { dg-do compile } */ + +#include +#include + +typedef struct { + int *a; + int *b; + int *c; +} mystruct; + +int main(int argc, char* argv[]) +{ + const int N = 1024; + const int S = 32; + mystruct *m = (mystruct *) calloc (S, sizeof (*m)); + int i, j; + + for (i = 0; i < S; i++) + { + m[i].a = (int *) malloc (N * sizeof (int)); + m[i].b = (int *) malloc (N * sizeof (int)); + m[i].c = (int *) malloc (N * sizeof (int)); + } + + for (j = 0; j < S; j++) + for (i = 0; i < N; i++) + { + m[j].a[i] = 0; + m[j].b[i] = 0; + m[j].c[i] = 0; + } + +#pragma acc enter data copyin(m[0:1]) + + for (int i = 0; i < 99; i++) + { + int j, k; + for (k = 0; k < S; k++) +#pragma acc parallel loop copy(m[k].a[0:N]) /* { dg-error "expected .\\\). before .\\\.. token" } */ + for (j = 0; j < N; j++) + m[k].a[j]++; + + for (k = 0; k < S; k++) +#pragma acc parallel loop copy(m[k].b[0:N], m[k].c[5:N-10]) /* { dg-error "expected .\\\). before .\\\.. token" } */ + /* { dg-error ".m. appears more than once in data clauses" "" { target c++ } .-1 } */ + for (j = 0; j < N; j++) + { + m[k].b[j]++; + if (j > 5 && j < N - 5) + m[k].c[j]++; + } + } + +#pragma acc exit data copyout(m[0:1]) + + for (j = 0; j < S; j++) + { + for (i = 0; i < N; i++) + { + if (m[j].a[i] != 99) + abort (); + if (m[j].b[i] != 99) + abort (); + if (i > 5 && i < N-5) + { + if (m[j].c[i] != 99) + abort (); + } + else + { + if (m[j].c[i] != 0) + abort (); + } + } + + free (m[j].a); + free (m[j].b); + free (m[j].c); + } + free (m); + + return 0; +} diff --git a/gcc/testsuite/c-c++-common/goacc/mdc-1.c b/gcc/testsuite/c-c++-common/goacc/mdc-1.c new file mode 100644 index 00000000000..6c6a81ea73a --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/mdc-1.c @@ -0,0 +1,55 @@ +/* Test OpenACC's support for manual deep copy, including the attach + and detach clauses. */ + +/* { dg-do compile { target int32 } } */ +/* { dg-additional-options "-fdump-tree-omplower" } */ + +void +t1 () +{ + struct foo { + int *a, *b, c, d, *e; + } s; + + int *a, *z; + +#pragma acc enter data copyin(s) + { +#pragma acc data copy(s.a[0:10]) copy(z[0:10]) + { + s.e = z; +#pragma acc parallel loop attach(s.e) + for (int i = 0; i < 10; i++) + s.a[i] = s.e[i]; + + + a = s.e; +#pragma acc enter data attach(a) +#pragma acc exit data detach(a) + } + +#pragma acc enter data copyin(a) +#pragma acc acc enter data attach(s.e) +#pragma acc exit data detach(s.e) + +#pragma acc data attach(s.e) + { + } +#pragma acc exit data delete(a) + +#pragma acc exit data detach(a) finalize +#pragma acc exit data detach(s.a) finalize + } +} + +/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.to:s .len: 32.." 1 "omplower" } } */ +/* { dg-final { scan-tree-dump-times "pragma omp target oacc_data map.tofrom:.z .len: 40.. map.struct:s .len: 1.. map.alloc:s.a .len: 8.. map.tofrom:._1 .len: 40.. map.attach:s.a .bias: 0.." 1 "omplower" } } */ +/* { dg-final { scan-tree-dump-times "pragma omp target oacc_parallel map.attach:s.e .bias: 8.. map.tofrom:s .len: 32" 1 "omplower" } } */ +/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.attach:a .bias: 8.." 1 "omplower" } } */ +/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.detach:a .bias: 8.." 1 "omplower" } } */ +/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.to:a .len: 8.." 1 "omplower" } } */ +/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.detach:s.e .bias: 8.." 1 "omplower" } } */ +/* { dg-final { scan-tree-dump-times "pragma omp target oacc_data map.attach:s.e .bias: 8.." 1 "omplower" } } */ +/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.release:a .len: 8.." 1 "omplower" } } */ +/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data finalize map.force_detach:a .bias: 8.." 1 "omplower" } } */ +/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data finalize map.force_detach:s.a .bias: 8.." 1 "omplower" } } */ diff --git a/gcc/testsuite/c-c++-common/goacc/mdc-2.c b/gcc/testsuite/c-c++-common/goacc/mdc-2.c new file mode 100644 index 00000000000..fae86671fc9 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/mdc-2.c @@ -0,0 +1,62 @@ +/* Test OpenACC's support for manual deep copy, including the attach + and detach clauses. */ + +void +t1 () +{ + struct foo { + int *a, *b, c, d, *e; + } s; + + int *a, *z, scalar, **y; + +#pragma acc enter data copyin(s) detach(z) /* { dg-error ".detach. is not valid for" } */ + { +#pragma acc data copy(s.a[0:10]) copy(z[0:10]) + { + s.e = z; +#pragma acc parallel loop attach(s.e) detach(s.b) /* { dg-error ".detach. is not valid for" } */ + for (int i = 0; i < 10; i++) + s.a[i] = s.e[i]; + + a = s.e; +#pragma acc enter data attach(a) detach(s.c) /* { dg-error ".detach. is not valid for" } */ +#pragma acc exit data detach(a) + } + +#pragma acc enter data attach(z[:5]) /* { dg-error "expected single pointer in .attach. clause" } */ +/* { dg-error "has no data movement clause" "" { target *-*-* } .-1 } */ +#pragma acc exit data detach(z[:5]) /* { dg-error "expected single pointer in .detach. clause" } */ +/* { dg-error "has no data movement clause" "" { target *-*-* } .-1 } */ +#pragma acc enter data attach(z[1:]) /* { dg-error "expected single pointer in .attach. clause" } */ +/* { dg-error "has no data movement clause" "" { target *-*-* } .-1 } */ +#pragma acc exit data detach(z[1:]) /* { dg-error "expected single pointer in .detach. clause" } */ +/* { dg-error "has no data movement clause" "" { target *-*-* } .-1 } */ +#pragma acc enter data attach(z[:]) /* { dg-error "expected single pointer in .attach. clause" } */ +/* { dg-error "has no data movement clause" "" { target *-*-* } .-1 } */ +#pragma acc exit data detach(z[:]) /* { dg-error "expected single pointer in .detach. clause" } */ +/* { dg-error "has no data movement clause" "" { target *-*-* } .-1 } */ +#pragma acc enter data attach(z[3]) /* { dg-error "expected pointer in .attach. clause" } */ +#pragma acc exit data detach(z[3]) /* { dg-error "expected pointer in .detach. clause" } */ + +#pragma acc acc enter data attach(s.e) +#pragma acc exit data detach(s.e) attach(z) /* { dg-error ".attach. is not valid for" } */ + +#pragma acc data attach(s.e) + { + } +#pragma acc exit data delete(a) attach(s.a) /* { dg-error ".attach. is not valid for" } */ + +#pragma acc enter data attach(scalar) /* { dg-error "expected pointer in .attach. clause" } */ +/* { dg-error "has no data movement clause" "" { target *-*-* } .-1 } */ +#pragma acc exit data detach(scalar) /* { dg-error "expected pointer in .detach. clause" } */ +/* { dg-error "has no data movement clause" "" { target *-*-* } .-1 } */ +#pragma acc enter data attach(s) /* { dg-error "expected pointer in .attach. clause" } */ +/* { dg-error "has no data movement clause" "" { target *-*-* } .-1 } */ +#pragma acc exit data detach(s) /* { dg-error "expected pointer in .detach. clause" } */ +/* { dg-error "has no data movement clause" "" { target *-*-* } .-1 } */ + } + +#pragma acc enter data attach(y[10]) +#pragma acc exit data detach(y[10]) +} diff --git a/gcc/testsuite/g++.dg/goacc/mdc.C b/gcc/testsuite/g++.dg/goacc/mdc.C new file mode 100644 index 00000000000..b3abab30423 --- /dev/null +++ b/gcc/testsuite/g++.dg/goacc/mdc.C @@ -0,0 +1,68 @@ +/* Test OpenACC's support for manual deep copy, including the attach + and detach clauses. */ + +void +t1 () +{ + struct foo { + int *a, *b, c, d, *e; + } s; + + struct foo& rs = s; + + int *a, *z, scalar, **y; + int* const &ra = a; + int* const &rz = z; + int& rscalar = scalar; + int** const &ry = y; + +#pragma acc enter data copyin(rs) detach(rz) /* { dg-error ".detach. is not valid for" } */ + { +#pragma acc data copy(rs.a[0:10]) copy(rz[0:10]) + { + s.e = z; +#pragma acc parallel loop attach(rs.e) detach(rs.b) /* { dg-error ".detach. is not valid for" } */ + for (int i = 0; i < 10; i++) + s.a[i] = s.e[i]; + + a = s.e; +#pragma acc enter data attach(ra) detach(rs.c) /* { dg-error ".detach. is not valid for" } */ +#pragma acc exit data detach(ra) + } + +#pragma acc enter data attach(rz[:5]) /* { dg-error "expected single pointer in .attach. clause" } */ +/* { dg-error "has no data movement clause" "" { target *-*-* } .-1 } */ +#pragma acc exit data detach(rz[:5]) /* { dg-error "expected single pointer in .detach. clause" } */ +/* { dg-error "has no data movement clause" "" { target *-*-* } .-1 } */ +#pragma acc enter data attach(rz[1:]) /* { dg-error "expected single pointer in .attach. clause" } */ +/* { dg-error "has no data movement clause" "" { target *-*-* } .-1 } */ +#pragma acc exit data detach(rz[1:]) /* { dg-error "expected single pointer in .detach. clause" } */ +/* { dg-error "has no data movement clause" "" { target *-*-* } .-1 } */ +#pragma acc enter data attach(rz[:]) /* { dg-error "expected single pointer in .attach. clause" } */ +/* { dg-error "has no data movement clause" "" { target *-*-* } .-1 } */ +#pragma acc exit data detach(rz[:]) /* { dg-error "expected single pointer in .detach. clause" } */ +/* { dg-error "has no data movement clause" "" { target *-*-* } .-1 } */ +#pragma acc enter data attach(rz[3]) /* { dg-error "expected pointer in .attach. clause" } */ +#pragma acc exit data detach(rz[3]) /* { dg-error "expected pointer in .detach. clause" } */ + +#pragma acc acc enter data attach(rs.e) +#pragma acc exit data detach(rs.e) attach(rz) /* { dg-error ".attach. is not valid for" } */ + +#pragma acc data attach(rs.e) + { + } +#pragma acc exit data delete(ra) attach(rs.a) /* { dg-error ".attach. is not valid for" } */ + +#pragma acc enter data attach(rscalar) /* { dg-error "expected pointer in .attach. clause" } */ +/* { dg-error "has no data movement clause" "" { target *-*-* } .-1 } */ +#pragma acc exit data detach(rscalar) /* { dg-error "expected pointer in .detach. clause" } */ +/* { dg-error "has no data movement clause" "" { target *-*-* } .-1 } */ +#pragma acc enter data attach(rs) /* { dg-error "expected pointer in .attach. clause" } */ +/* { dg-error "has no data movement clause" "" { target *-*-* } .-1 } */ +#pragma acc exit data detach(rs) /* { dg-error "expected pointer in .detach. clause" } */ +/* { dg-error "has no data movement clause" "" { target *-*-* } .-1 } */ + } + +#pragma acc enter data attach(ry[10]) +#pragma acc exit data detach(ry[10]) +} diff --git a/gcc/testsuite/gfortran.dg/goacc/data-clauses.f95 b/gcc/testsuite/gfortran.dg/goacc/data-clauses.f95 index b94214e8b63..1a4a6719987 100644 --- a/gcc/testsuite/gfortran.dg/goacc/data-clauses.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/data-clauses.f95 @@ -39,9 +39,9 @@ contains !$acc end data - !$acc parallel copy (tip) ! { dg-error "POINTER" } + !$acc parallel copy (tip) !$acc end parallel - !$acc parallel copy (tia) ! { dg-error "ALLOCATABLE" } + !$acc parallel copy (tia) !$acc end parallel !$acc parallel deviceptr (i) copy (i) ! { dg-error "multiple clauses" } !$acc end parallel @@ -54,9 +54,9 @@ contains !$acc end data - !$acc parallel copyin (tip) ! { dg-error "POINTER" } + !$acc parallel copyin (tip) !$acc end parallel - !$acc parallel copyin (tia) ! { dg-error "ALLOCATABLE" } + !$acc parallel copyin (tia) !$acc end parallel !$acc parallel deviceptr (i) copyin (i) ! { dg-error "multiple clauses" } !$acc end parallel @@ -71,9 +71,9 @@ contains !$acc end data - !$acc parallel copyout (tip) ! { dg-error "POINTER" } + !$acc parallel copyout (tip) !$acc end parallel - !$acc parallel copyout (tia) ! { dg-error "ALLOCATABLE" } + !$acc parallel copyout (tia) !$acc end parallel !$acc parallel deviceptr (i) copyout (i) ! { dg-error "multiple clauses" } !$acc end parallel @@ -90,9 +90,9 @@ contains !$acc end data - !$acc parallel create (tip) ! { dg-error "POINTER" } + !$acc parallel create (tip) !$acc end parallel - !$acc parallel create (tia) ! { dg-error "ALLOCATABLE" } + !$acc parallel create (tia) !$acc end parallel !$acc parallel deviceptr (i) create (i) ! { dg-error "multiple clauses" } !$acc end parallel @@ -111,9 +111,9 @@ contains !$acc end data - !$acc parallel present (tip) ! { dg-error "POINTER" } + !$acc parallel present (tip) !$acc end parallel - !$acc parallel present (tia) ! { dg-error "ALLOCATABLE" } + !$acc parallel present (tia) !$acc end parallel !$acc parallel deviceptr (i) present (i) ! { dg-error "multiple clauses" } !$acc end parallel @@ -144,9 +144,9 @@ contains !$acc end parallel - !$acc parallel present_or_copy (tip) ! { dg-error "POINTER" } + !$acc parallel present_or_copy (tip) !$acc end parallel - !$acc parallel present_or_copy (tia) ! { dg-error "ALLOCATABLE" } + !$acc parallel present_or_copy (tia) !$acc end parallel !$acc parallel deviceptr (i) present_or_copy (i) ! { dg-error "multiple clauses" } !$acc end parallel @@ -169,9 +169,9 @@ contains !$acc end data - !$acc parallel present_or_copyin (tip) ! { dg-error "POINTER" } + !$acc parallel present_or_copyin (tip) !$acc end parallel - !$acc parallel present_or_copyin (tia) ! { dg-error "ALLOCATABLE" } + !$acc parallel present_or_copyin (tia) !$acc end parallel !$acc parallel deviceptr (i) present_or_copyin (i) ! { dg-error "multiple clauses" } !$acc end parallel @@ -196,9 +196,9 @@ contains !$acc end data - !$acc parallel present_or_copyout (tip) ! { dg-error "POINTER" } + !$acc parallel present_or_copyout (tip) !$acc end parallel - !$acc parallel present_or_copyout (tia) ! { dg-error "ALLOCATABLE" } + !$acc parallel present_or_copyout (tia) !$acc end parallel !$acc parallel deviceptr (i) present_or_copyout (i) ! { dg-error "multiple clauses" } !$acc end parallel @@ -225,9 +225,9 @@ contains !$acc end data - !$acc parallel present_or_create (tip) ! { dg-error "POINTER" } + !$acc parallel present_or_create (tip) !$acc end parallel - !$acc parallel present_or_create (tia) ! { dg-error "ALLOCATABLE" } + !$acc parallel present_or_create (tia) !$acc end parallel !$acc parallel deviceptr (i) present_or_create (i) ! { dg-error "multiple clauses" } !$acc end parallel @@ -256,4 +256,4 @@ contains !$acc end data end subroutine foo -end module test \ No newline at end of file +end module test diff --git a/gcc/testsuite/gfortran.dg/goacc/derived-types-2.f90 b/gcc/testsuite/gfortran.dg/goacc/derived-types-2.f90 new file mode 100644 index 00000000000..d01583fac89 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/goacc/derived-types-2.f90 @@ -0,0 +1,14 @@ +module bar + type :: type1 + real(8), pointer, public :: p(:) => null() + end type + type :: type2 + class(type1), pointer :: p => null() + end type +end module + +subroutine foo (var) + use bar + type(type2), intent(inout) :: var + !$acc enter data create(var%p%p) +end subroutine diff --git a/gcc/testsuite/gfortran.dg/goacc/derived-types.f90 b/gcc/testsuite/gfortran.dg/goacc/derived-types.f90 new file mode 100644 index 00000000000..5fb29816c42 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/goacc/derived-types.f90 @@ -0,0 +1,77 @@ +! Test ACC UPDATE with derived types. + +module dt + integer, parameter :: n = 10 + type inner + integer :: d(n) + end type inner + type dtype + integer(8) :: a, b, c(n) + type(inner) :: in + end type dtype +end module dt + +program derived_acc + use dt + + implicit none + type(dtype):: var + integer i + !$acc declare create(var) + !$acc declare pcopy(var%a) ! { dg-error "Syntax error in OpenMP" } + + !$acc update host(var) + !$acc update host(var%a) + !$acc update device(var) + !$acc update device(var%a) + !$acc update self(var) + !$acc update self(var%a) + + !$acc enter data copyin(var) + !$acc enter data copyin(var%a) + + !$acc exit data copyout(var) + !$acc exit data copyout(var%a) + + !$acc data copy(var) + !$acc end data + + !$acc data copyout(var%a) + !$acc end data + + !$acc parallel loop pcopyout(var) + do i = 1, 10 + end do + !$acc end parallel loop + + !$acc parallel loop copyout(var%a) + do i = 1, 10 + end do + !$acc end parallel loop + + !$acc parallel pcopy(var) + !$acc end parallel + + !$acc parallel pcopy(var%a) + do i = 1, 10 + end do + !$acc end parallel + + !$acc kernels pcopyin(var) + !$acc end kernels + + !$acc kernels pcopy(var%a) + do i = 1, 10 + end do + !$acc end kernels + + !$acc kernels loop pcopyin(var) + do i = 1, 10 + end do + !$acc end kernels loop + + !$acc kernels loop pcopy(var%a) + do i = 1, 10 + end do + !$acc end kernels loop +end program derived_acc diff --git a/gcc/testsuite/gfortran.dg/goacc/enter-exit-data.f95 b/gcc/testsuite/gfortran.dg/goacc/enter-exit-data.f95 index a414df8d439..c2a49796318 100644 --- a/gcc/testsuite/gfortran.dg/goacc/enter-exit-data.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/enter-exit-data.f95 @@ -44,14 +44,14 @@ contains !$acc enter data wait (i, 1) !$acc enter data wait (a) ! { dg-error "INTEGER" } !$acc enter data wait (b(5:6)) ! { dg-error "INTEGER" } - !$acc enter data copyin (tip) ! { dg-error "POINTER" } - !$acc enter data copyin (tia) ! { dg-error "ALLOCATABLE" } - !$acc enter data create (tip) ! { dg-error "POINTER" } - !$acc enter data create (tia) ! { dg-error "ALLOCATABLE" } - !$acc enter data present_or_copyin (tip) ! { dg-error "POINTER" } - !$acc enter data present_or_copyin (tia) ! { dg-error "ALLOCATABLE" } - !$acc enter data present_or_create (tip) ! { dg-error "POINTER" } - !$acc enter data present_or_create (tia) ! { dg-error "ALLOCATABLE" } + !$acc enter data copyin (tip) + !$acc enter data copyin (tia) + !$acc enter data create (tip) + !$acc enter data create (tia) + !$acc enter data present_or_copyin (tip) + !$acc enter data present_or_copyin (tia) + !$acc enter data present_or_create (tip) + !$acc enter data present_or_create (tia) !$acc enter data copyin (i) create (i) ! { dg-error "multiple clauses" } !$acc enter data copyin (i) present_or_copyin (i) ! { dg-error "multiple clauses" } !$acc enter data create (i) present_or_copyin (i) ! { dg-error "multiple clauses" } @@ -79,10 +79,10 @@ contains !$acc exit data wait (i, 1) !$acc exit data wait (a) ! { dg-error "INTEGER" } !$acc exit data wait (b(5:6)) ! { dg-error "INTEGER" } - !$acc exit data copyout (tip) ! { dg-error "POINTER" } - !$acc exit data copyout (tia) ! { dg-error "ALLOCATABLE" } - !$acc exit data delete (tip) ! { dg-error "POINTER" } - !$acc exit data delete (tia) ! { dg-error "ALLOCATABLE" } + !$acc exit data copyout (tip) + !$acc exit data copyout (tia) + !$acc exit data delete (tip) + !$acc exit data delete (tia) !$acc exit data copyout (i) delete (i) ! { dg-error "multiple clauses" } !$acc exit data finalize !$acc exit data finalize copyout (i) diff --git a/gcc/tree-pretty-print.c b/gcc/tree-pretty-print.c index 53b3f55a3e6..df617d47899 100644 --- a/gcc/tree-pretty-print.c +++ b/gcc/tree-pretty-print.c @@ -849,6 +849,18 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags) case GOMP_MAP_LINK: pp_string (pp, "link"); break; + case GOMP_MAP_ATTACH: + pp_string (pp, "attach"); + break; + case GOMP_MAP_DETACH: + pp_string (pp, "detach"); + break; + case GOMP_MAP_FORCE_DETACH: + pp_string (pp, "force_detach"); + break; + case GOMP_MAP_ATTACH_DETACH: + pp_string (pp, "attach_detach"); + break; default: gcc_unreachable (); } @@ -870,6 +882,12 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags) case GOMP_MAP_TO_PSET: pp_string (pp, " [pointer set, len: "); break; + case GOMP_MAP_ATTACH: + case GOMP_MAP_DETACH: + case GOMP_MAP_FORCE_DETACH: + case GOMP_MAP_ATTACH_DETACH: + pp_string (pp, " [bias: "); + break; default: pp_string (pp, " [len: "); break; diff --git a/include/gomp-constants.h b/include/gomp-constants.h index 82e9094c934..6aa7691a119 100644 --- a/include/gomp-constants.h +++ b/include/gomp-constants.h @@ -40,8 +40,11 @@ #define GOMP_MAP_FLAG_SPECIAL_0 (1 << 2) #define GOMP_MAP_FLAG_SPECIAL_1 (1 << 3) #define GOMP_MAP_FLAG_SPECIAL_2 (1 << 4) +#define GOMP_MAP_FLAG_SPECIAL_4 (1 << 6) #define GOMP_MAP_FLAG_SPECIAL (GOMP_MAP_FLAG_SPECIAL_1 \ | GOMP_MAP_FLAG_SPECIAL_0) +#define GOMP_MAP_DEEP_COPY (GOMP_MAP_FLAG_SPECIAL_4 \ + | GOMP_MAP_FLAG_SPECIAL_2) /* Flag to force a specific behavior (or else, trigger a run-time error). */ #define GOMP_MAP_FLAG_FORCE (1 << 7) @@ -127,12 +130,23 @@ enum gomp_map_kind /* Decrement usage count and deallocate if zero. */ GOMP_MAP_RELEASE = (GOMP_MAP_FLAG_SPECIAL_2 | GOMP_MAP_DELETE), + /* In OpenACC, attach a pointer to a mapped struct field. */ + GOMP_MAP_ATTACH = (GOMP_MAP_DEEP_COPY | 0), + /* In OpenACC, detach a pointer to a mapped struct field. */ + GOMP_MAP_DETACH = (GOMP_MAP_DEEP_COPY | 1), + /* In OpenACC, detach a pointer to a mapped struct field. */ + GOMP_MAP_FORCE_DETACH = (GOMP_MAP_DEEP_COPY + | GOMP_MAP_FLAG_FORCE | 1), /* Internal to GCC, not used in libgomp. */ /* Do not map, but pointer assign a pointer instead. */ GOMP_MAP_FIRSTPRIVATE_POINTER = (GOMP_MAP_LAST | 1), /* Do not map, but pointer assign a reference instead. */ - GOMP_MAP_FIRSTPRIVATE_REFERENCE = (GOMP_MAP_LAST | 2) + GOMP_MAP_FIRSTPRIVATE_REFERENCE = (GOMP_MAP_LAST | 2), + /* An attach or detach operation. Rewritten to the appropriate type during + gimplification, depending on directive (i.e. "enter data" or + parallel/kernels region vs. "exit data"). */ + GOMP_MAP_ATTACH_DETACH = (GOMP_MAP_LAST | 3) }; #define GOMP_MAP_COPY_TO_P(X) \ diff --git a/libgomp/libgomp.h b/libgomp/libgomp.h index 6870ee34224..9953fabc586 100644 --- a/libgomp/libgomp.h +++ b/libgomp/libgomp.h @@ -867,6 +867,8 @@ struct target_var_desc { bool copy_from; /* True if data always should be copied from device to host at the end. */ bool always_copy_from; + /* True if variable should be detached at end of region. */ + bool do_detach; /* Relative offset against key host_start. */ uintptr_t offset; /* Actual length. */ @@ -943,6 +945,9 @@ struct splay_tree_key_s { uintptr_t refcount_chk; #endif union { + /* For a block with attached pointers, the attachment counters for each. + Only used for OpenACC. */ + uintptr_t *attach_count; /* Pointer to the original mapping of "omp declare target link" object. Only used for OpenMP. */ splay_tree_key link_key; @@ -1096,6 +1101,13 @@ extern void gomp_copy_dev2host (struct gomp_device_descr *, struct goacc_asyncqueue *, void *, const void *, size_t); extern uintptr_t gomp_map_val (struct target_mem_desc *, void **, size_t); +extern void gomp_attach_pointer (struct gomp_device_descr *, + struct goacc_asyncqueue *, splay_tree, + splay_tree_key, uintptr_t, size_t, + struct gomp_coalesce_buf *); +extern void gomp_detach_pointer (struct gomp_device_descr *, + struct goacc_asyncqueue *, splay_tree_key, + uintptr_t, bool, struct gomp_coalesce_buf *); #ifdef RC_CHECKING extern void dump_tgt (const char *, struct target_mem_desc *); diff --git a/libgomp/libgomp.map b/libgomp/libgomp.map index c79430f8d8d..63276f7d29b 100644 --- a/libgomp/libgomp.map +++ b/libgomp/libgomp.map @@ -484,6 +484,16 @@ OACC_2.5.1 { acc_register_library; } OACC_2.5; +OACC_2.6 { + global: + acc_attach; + acc_attach_async; + acc_detach; + acc_detach_async; + acc_detach_finalize; + acc_detach_finalize_async; +} OACC_2.5.1; + GOACC_2.0 { global: GOACC_data_end; diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c index 25084b71a2d..c88fe9732d0 100644 --- a/libgomp/oacc-mem.c +++ b/libgomp/oacc-mem.c @@ -875,9 +875,13 @@ gomp_acc_remove_pointer (struct gomp_device_descr *acc_dev, void **hostaddrs, case GOMP_MAP_POINTER: case GOMP_MAP_DELETE: case GOMP_MAP_RELEASE: + case GOMP_MAP_DETACH: + case GOMP_MAP_FORCE_DETACH: cur_node.host_start = (uintptr_t) hostaddrs[i]; cur_node.host_end = cur_node.host_start - + (kind == GOMP_MAP_POINTER + + ((kind == GOMP_MAP_DETACH + || kind == GOMP_MAP_FORCE_DETACH + || kind == GOMP_MAP_POINTER) ? sizeof (void *) : sizes[i]); n = splay_tree_lookup (&acc_dev->mem_map, &cur_node); @@ -926,3 +930,79 @@ gomp_acc_remove_pointer (struct gomp_device_descr *acc_dev, void **hostaddrs, gomp_mutex_unlock (&acc_dev->lock); } + +void +acc_attach_async (void **hostaddr, int async) +{ + struct goacc_thread *thr = goacc_thread (); + struct gomp_device_descr *acc_dev = thr->dev; + goacc_aq aq = get_goacc_asyncqueue (async); + + struct splay_tree_key_s cur_node; + splay_tree_key n; + + if (thr->dev->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM) + return; + + cur_node.host_start = (uintptr_t) hostaddr; + cur_node.host_end = cur_node.host_start + sizeof (void *); + n = splay_tree_lookup (&acc_dev->mem_map, &cur_node); + + if (n == NULL) + gomp_fatal ("struct not mapped for acc_attach"); + + gomp_attach_pointer (acc_dev, aq, &acc_dev->mem_map, n, (uintptr_t) hostaddr, + 0, NULL); +} + +void +acc_attach (void **hostaddr) +{ + acc_attach_async (hostaddr, acc_async_sync); +} + +static void +goacc_detach_internal (void **hostaddr, int async, bool finalize) +{ + struct goacc_thread *thr = goacc_thread (); + struct gomp_device_descr *acc_dev = thr->dev; + struct splay_tree_key_s cur_node; + splay_tree_key n; + struct goacc_asyncqueue *aq = get_goacc_asyncqueue (async); + + if (thr->dev->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM) + return; + + cur_node.host_start = (uintptr_t) hostaddr; + cur_node.host_end = cur_node.host_start + sizeof (void *); + n = splay_tree_lookup (&acc_dev->mem_map, &cur_node); + + if (n == NULL) + gomp_fatal ("struct not mapped for acc_detach"); + + gomp_detach_pointer (acc_dev, aq, n, (uintptr_t) hostaddr, finalize, NULL); +} + +void +acc_detach (void **hostaddr) +{ + goacc_detach_internal (hostaddr, acc_async_sync, false); +} + +void +acc_detach_async (void **hostaddr, int async) +{ + goacc_detach_internal (hostaddr, async, false); +} + +void +acc_detach_finalize (void **hostaddr) +{ + goacc_detach_internal (hostaddr, acc_async_sync, true); +} + +void +acc_detach_finalize_async (void **hostaddr, int async) +{ + goacc_detach_internal (hostaddr, async, true); +} diff --git a/libgomp/oacc-parallel.c b/libgomp/oacc-parallel.c index 02472c56270..69b9f379476 100644 --- a/libgomp/oacc-parallel.c +++ b/libgomp/oacc-parallel.c @@ -71,7 +71,10 @@ find_pointer (int pos, size_t mapnum, unsigned short *kinds) { unsigned char kind1 = kinds[pos + 1] & 0xff; if (kind1 == GOMP_MAP_POINTER - || kind1 == GOMP_MAP_ALWAYS_POINTER) + || kind1 == GOMP_MAP_ALWAYS_POINTER + || kind1 == GOMP_MAP_ATTACH + || kind1 == GOMP_MAP_DETACH + || kind1 == GOMP_MAP_FORCE_DETACH) return 2; else if (kind1 == GOMP_MAP_TO_PSET) return 3; @@ -643,6 +646,10 @@ GOACC_enter_exit_data (int flags_m, size_t mapnum, if (mapnum > 0) { unsigned char kind = kinds[0] & 0xff; + + if (kind == GOMP_MAP_STRUCT || kind == GOMP_MAP_FORCE_PRESENT) + kind = kinds[1] & 0xff; + if (kind == GOMP_MAP_DELETE || kind == GOMP_MAP_FORCE_FROM) finalize = true; @@ -653,11 +660,14 @@ GOACC_enter_exit_data (int flags_m, size_t mapnum, { unsigned char kind = kinds[i] & 0xff; - if (kind == GOMP_MAP_POINTER || kind == GOMP_MAP_TO_PSET) + if (kind == GOMP_MAP_POINTER + || kind == GOMP_MAP_TO_PSET + || kind == GOMP_MAP_STRUCT) continue; if (kind == GOMP_MAP_FORCE_ALLOC || kind == GOMP_MAP_FORCE_PRESENT + || kind == GOMP_MAP_ATTACH || kind == GOMP_MAP_FORCE_TO || kind == GOMP_MAP_TO || kind == GOMP_MAP_ALLOC) @@ -668,6 +678,8 @@ GOACC_enter_exit_data (int flags_m, size_t mapnum, if (kind == GOMP_MAP_RELEASE || kind == GOMP_MAP_DELETE + || kind == GOMP_MAP_DETACH + || kind == GOMP_MAP_FORCE_DETACH || kind == GOMP_MAP_FROM || kind == GOMP_MAP_FORCE_FROM) break; @@ -777,6 +789,19 @@ GOACC_enter_exit_data (int flags_m, size_t mapnum, case GOMP_MAP_FORCE_TO: acc_copyin_async (hostaddrs[i], sizes[i], async); break; + case GOMP_MAP_STRUCT: + { + int elems = sizes[i]; + goacc_aq aq = get_goacc_asyncqueue (async); + gomp_map_vars_async (acc_dev, aq, elems + 1, &hostaddrs[i], + NULL, &sizes[i], &kinds[i], true, + GOMP_MAP_VARS_OPENACC_ENTER_DATA); + i += elems; + } + break; + case GOMP_MAP_ATTACH: + case GOMP_MAP_FORCE_PRESENT: + break; default: gomp_fatal (">>>> GOACC_enter_exit_data UNHANDLED kind 0x%.2x", kind); @@ -800,48 +825,140 @@ GOACC_enter_exit_data (int flags_m, size_t mapnum, i += pointer - 1; } } + + /* This loop only handles explicit "attach" clauses that are not an + implicit part of a copy{,in,out}, etc. mapping. */ + for (i = 0; i < mapnum; i++) + { + unsigned char kind = kinds[i] & 0xff; + + /* Scan for pointers and PSETs. */ + int pointer = find_pointer (i, mapnum, kinds); + + if (!pointer) + { + if (kind == GOMP_MAP_ATTACH) + acc_attach (hostaddrs[i]); + else if (kind == GOMP_MAP_STRUCT) + i += sizes[i]; + } + else + i += pointer - 1; + } } else - for (i = 0; i < mapnum; ++i) - { - unsigned char kind = kinds[i] & 0xff; + { + /* Handle "detach" before copyback/deletion of mapped data. */ + for (i = 0; i < mapnum; i++) + { + unsigned char kind = kinds[i] & 0xff; - int pointer = find_pointer (i, mapnum, kinds); + int pointer = find_pointer (i, mapnum, kinds); - if (!pointer) - { - switch (kind) - { - case GOMP_MAP_RELEASE: - case GOMP_MAP_DELETE: - if (acc_is_present (hostaddrs[i], sizes[i])) + if (!pointer) + { + if (kind == GOMP_MAP_DETACH) + acc_detach (hostaddrs[i]); + else if (kind == GOMP_MAP_FORCE_DETACH) + acc_detach_finalize (hostaddrs[i]); + else if (kind == GOMP_MAP_STRUCT) + i += sizes[i]; + } + else + { + unsigned char kind2 = kinds[i + pointer - 1] & 0xff; + + if (kind2 == GOMP_MAP_DETACH) + acc_detach (hostaddrs[i + pointer - 1]); + else if (kind2 == GOMP_MAP_FORCE_DETACH) + acc_detach_finalize (hostaddrs[i + pointer - 1]); + + i += pointer - 1; + } + } + + for (i = 0; i < mapnum; ++i) + { + unsigned char kind = kinds[i] & 0xff; + + int pointer = find_pointer (i, mapnum, kinds); + + if (!pointer) + { + switch (kind) + { + case GOMP_MAP_RELEASE: + case GOMP_MAP_DELETE: + if (acc_is_present (hostaddrs[i], sizes[i])) + { + if (finalize) + acc_delete_finalize_async (hostaddrs[i], sizes[i], + async); + else + acc_delete_async (hostaddrs[i], sizes[i], async); + } + break; + case GOMP_MAP_DETACH: + case GOMP_MAP_FORCE_DETACH: + case GOMP_MAP_FORCE_PRESENT: + break; + case GOMP_MAP_FROM: + case GOMP_MAP_FORCE_FROM: + if (finalize) + acc_copyout_finalize_async (hostaddrs[i], sizes[i], async); + else + acc_copyout_async (hostaddrs[i], sizes[i], async); + break; + case GOMP_MAP_STRUCT: { - if (finalize) - acc_delete_finalize_async (hostaddrs[i], sizes[i], async); - else - acc_delete_async (hostaddrs[i], sizes[i], async); + int elems = sizes[i]; + goacc_aq aq = get_goacc_asyncqueue (async); + for (int j = 1; j <= elems; j++) + { + struct splay_tree_key_s k; + k.host_start = (uintptr_t) hostaddrs[i + j]; + k.host_end = k.host_start + sizes[i + j]; + splay_tree_key str; + gomp_mutex_lock (&acc_dev->lock); + str = splay_tree_lookup (&acc_dev->mem_map, &k); + gomp_mutex_unlock (&acc_dev->lock); + if (str) + { + assert (str->virtual_refcount + != VREFCOUNT_LINK_KEY); + if (finalize) + { + str->refcount -= str->virtual_refcount; + str->virtual_refcount = 0; + } + if (str->virtual_refcount > 0) + { + str->refcount--; + str->virtual_refcount--; + } + else if (str->refcount > 0) + str->refcount--; + if (str->refcount == 0) + gomp_remove_var_async (acc_dev, str, aq); + } + } + i += elems; } - break; - case GOMP_MAP_FROM: - case GOMP_MAP_FORCE_FROM: - if (finalize) - acc_copyout_finalize_async (hostaddrs[i], sizes[i], async); - else - acc_copyout_async (hostaddrs[i], sizes[i], async); - break; - default: - gomp_fatal (">>>> GOACC_enter_exit_data UNHANDLED kind 0x%.2x", - kind); - break; - } - } - else - { - gomp_acc_remove_pointer (acc_dev, &hostaddrs[i], &sizes[i], - &kinds[i], async, finalize, pointer); - i += pointer - 1; - } - } + break; + default: + gomp_fatal (">>>> GOACC_enter_exit_data UNHANDLED kind " + "0x%.2x", kind); + break; + } + } + else + { + gomp_acc_remove_pointer (acc_dev, &hostaddrs[i], &sizes[i], + &kinds[i], async, finalize, pointer); + i += pointer - 1; + } + } + } out_prof: if (profiling_p) diff --git a/libgomp/openacc.h b/libgomp/openacc.h index 1bbe6c90e7f..afe31cfd99b 100644 --- a/libgomp/openacc.h +++ b/libgomp/openacc.h @@ -108,12 +108,18 @@ void *acc_hostptr (void *) __GOACC_NOTHROW; int acc_is_present (void *, size_t) __GOACC_NOTHROW; void acc_memcpy_to_device (void *, void *, size_t) __GOACC_NOTHROW; void acc_memcpy_from_device (void *, void *, size_t) __GOACC_NOTHROW; +void acc_attach (void **) __GOACC_NOTHROW; +void acc_attach_async (void **, int) __GOACC_NOTHROW; +void acc_detach (void **) __GOACC_NOTHROW; +void acc_detach_async (void **, int) __GOACC_NOTHROW; /* Finalize versions of copyout/delete functions, specified in OpenACC 2.5. */ void acc_copyout_finalize (void *, size_t) __GOACC_NOTHROW; void acc_copyout_finalize_async (void *, size_t, int) __GOACC_NOTHROW; void acc_delete_finalize (void *, size_t) __GOACC_NOTHROW; void acc_delete_finalize_async (void *, size_t, int) __GOACC_NOTHROW; +void acc_detach_finalize (void **) __GOACC_NOTHROW; +void acc_detach_finalize_async (void **, int) __GOACC_NOTHROW; /* Async functions, specified in OpenACC 2.5. */ void acc_copyin_async (void *, size_t, int) __GOACC_NOTHROW; diff --git a/libgomp/target.c b/libgomp/target.c index aa7a1df1e46..eb28591ee4f 100644 --- a/libgomp/target.c +++ b/libgomp/target.c @@ -381,8 +381,12 @@ dump_tgt (const char *where, struct target_mem_desc *tgt) fprintf (stderr, " key.u.link_key=%p\n", (void*) tgt->list[i].key->u.link_key); else - fprintf (stderr, " key.virtual_refcount=%d\n", - (int) tgt->list[i].key->virtual_refcount); + { + fprintf (stderr, " key.virtual_refcount=%d\n", + (int) tgt->list[i].key->virtual_refcount); + fprintf (stderr, " key.u.attach_count=%p\n", + (void*) tgt->list[i].key->u.attach_count); + } } } fprintf (stderr, "\n"); @@ -537,6 +541,7 @@ gomp_map_vars_existing (struct gomp_device_descr *devicep, tgt_var->key = oldn; tgt_var->copy_from = GOMP_MAP_COPY_FROM_P (kind); tgt_var->always_copy_from = GOMP_MAP_ALWAYS_FROM_P (kind); + tgt_var->do_detach = kind == GOMP_MAP_ATTACH; tgt_var->offset = newn->host_start - oldn->host_start; tgt_var->length = newn->host_end - newn->host_start; @@ -670,6 +675,135 @@ gomp_map_fields_existing (struct target_mem_desc *tgt, (void *) cur_node.host_end); } +void +gomp_attach_pointer (struct gomp_device_descr *devicep, + struct goacc_asyncqueue *aq, splay_tree mem_map, + splay_tree_key n, uintptr_t attach_to, size_t bias, + struct gomp_coalesce_buf *cbufp) +{ + struct splay_tree_key_s s; + size_t size, idx; + + if (n == NULL) + { + gomp_mutex_unlock (&devicep->lock); + gomp_fatal ("enclosing struct not mapped for attach"); + } + + size = (n->host_end - n->host_start + sizeof (void *) - 1) / sizeof (void *); + /* We might have a pointer in a packed struct: however we cannot have more + than one such pointer in each pointer-sized portion of the struct, so + this is safe. */ + idx = (attach_to - n->host_start) / sizeof (void *); + + assert (n->virtual_refcount != VREFCOUNT_LINK_KEY); + + if (!n->u.attach_count) + n->u.attach_count + = gomp_malloc_cleared (sizeof (*n->u.attach_count) * size); + + if (n->u.attach_count[idx] < UINTPTR_MAX) + n->u.attach_count[idx]++; + else + { + gomp_mutex_unlock (&devicep->lock); + gomp_fatal ("attach count overflow"); + } + + if (n->u.attach_count[idx] == 1) + { + uintptr_t devptr = n->tgt->tgt_start + n->tgt_offset + attach_to + - n->host_start; + uintptr_t target = (uintptr_t) *(void **) attach_to; + splay_tree_key tn; + uintptr_t data; + + if ((void *) target == NULL) + { + gomp_mutex_unlock (&devicep->lock); + gomp_fatal ("attempt to attach null pointer"); + } + + s.host_start = target + bias; + s.host_end = s.host_start + 1; + tn = splay_tree_lookup (mem_map, &s); + + if (!tn) + { + gomp_mutex_unlock (&devicep->lock); + gomp_fatal ("pointer target not mapped for attach"); + } + + data = tn->tgt->tgt_start + tn->tgt_offset + target - tn->host_start; + + gomp_debug (1, + "%s: attaching host %p, target %p (struct base %p) to %p\n", + __FUNCTION__, (void *) attach_to, (void *) devptr, + (void *) (n->tgt->tgt_start + n->tgt_offset), (void *) data); + + gomp_copy_host2dev (devicep, aq, (void *) devptr, (void *) &data, + sizeof (void *), cbufp); + } + else + gomp_debug (1, "%s: attach count for %p -> %u\n", __FUNCTION__, + (void *) attach_to, (int) n->u.attach_count[idx]); +} + +void +gomp_detach_pointer (struct gomp_device_descr *devicep, + struct goacc_asyncqueue *aq, splay_tree_key n, + uintptr_t detach_from, bool finalize, + struct gomp_coalesce_buf *cbufp) +{ + size_t idx; + + if (n == NULL) + { + gomp_mutex_unlock (&devicep->lock); + gomp_fatal ("enclosing struct not mapped for detach"); + } + + idx = (detach_from - n->host_start) / sizeof (void *); + + assert (n->virtual_refcount != VREFCOUNT_LINK_KEY); + + if (!n->u.attach_count) + { + gomp_mutex_unlock (&devicep->lock); + gomp_fatal ("no attachment counters for struct"); + } + + if (finalize) + n->u.attach_count[idx] = 1; + + if (n->u.attach_count[idx] == 0) + { + gomp_mutex_unlock (&devicep->lock); + gomp_fatal ("attach count underflow"); + } + else + n->u.attach_count[idx]--; + + if (n->u.attach_count[idx] == 0) + { + uintptr_t devptr = n->tgt->tgt_start + n->tgt_offset + detach_from + - n->host_start; + uintptr_t target = (uintptr_t) *(void **) detach_from; + + gomp_debug (1, + "%s: detaching host %p, target %p (struct base %p) to %p\n", + __FUNCTION__, (void *) detach_from, (void *) devptr, + (void *) (n->tgt->tgt_start + n->tgt_offset), + (void *) target); + + gomp_copy_host2dev (devicep, aq, (void *) devptr, (void *) &target, + sizeof (void *), cbufp); + } + else + gomp_debug (1, "%s: attach count for %p -> %u\n", __FUNCTION__, + (void *) detach_from, (int) n->u.attach_count[idx]); +} + uintptr_t gomp_map_val (struct target_mem_desc *tgt, void **hostaddrs, size_t i) { @@ -822,8 +956,15 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep, has_firstprivate = true; continue; } + else if ((kind & typemask) == GOMP_MAP_ATTACH) + { + tgt->list[i].key = NULL; + has_firstprivate = true; + continue; + } cur_node.host_start = (uintptr_t) hostaddrs[i]; - if (!GOMP_MAP_POINTER_P (kind & typemask)) + if (!GOMP_MAP_POINTER_P (kind & typemask) + && (kind & typemask) != GOMP_MAP_ATTACH) cur_node.host_end = cur_node.host_start + sizes[i]; else cur_node.host_end = cur_node.host_start + sizeof (void *); @@ -1047,6 +1188,32 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep, cur_node.tgt_offset = n->tgt->tgt_start + n->tgt_offset + cur_node.host_start - n->host_start; continue; + case GOMP_MAP_ATTACH: + { + cur_node.host_start = (uintptr_t) hostaddrs[i]; + cur_node.host_end = cur_node.host_start + sizeof (void *); + splay_tree_key n = splay_tree_lookup (mem_map, &cur_node); + if (n != NULL) + { + tgt->list[i].key = n; + tgt->list[i].offset = cur_node.host_start - n->host_start; + tgt->list[i].length = n->host_end - n->host_start; + tgt->list[i].copy_from = false; + tgt->list[i].always_copy_from = false; + tgt->list[i].do_detach + = (pragma_kind != GOMP_MAP_VARS_OPENACC_ENTER_DATA); + n->refcount++; + } + else + { + gomp_mutex_unlock (&devicep->lock); + gomp_fatal ("outer struct not mapped for attach"); + } + gomp_attach_pointer (devicep, aq, mem_map, n, + (uintptr_t) hostaddrs[i], sizes[i], + cbufp); + continue; + } default: break; } @@ -1090,10 +1257,12 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep, tgt->list[i].copy_from = GOMP_MAP_COPY_FROM_P (kind & typemask); tgt->list[i].always_copy_from = GOMP_MAP_ALWAYS_FROM_P (kind & typemask); + tgt->list[i].do_detach = false; tgt->list[i].offset = 0; tgt->list[i].length = k->host_end - k->host_start; k->refcount = 1; k->virtual_refcount = 0; + k->u.attach_count = NULL; tgt->refcount++; array->left = NULL; array->right = NULL; @@ -1144,6 +1313,7 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep, tgt->list[j].key = k; tgt->list[j].copy_from = false; tgt->list[j].always_copy_from = false; + tgt->list[j].do_detach = false; if (k->refcount != REFCOUNT_INFINITY) k->refcount++; gomp_map_pointer (tgt, aq, @@ -1319,6 +1489,8 @@ gomp_remove_var_internal (struct gomp_device_descr *devicep, splay_tree_key k, if (k->u.link_key) splay_tree_insert (&devicep->mem_map, (splay_tree_node) k->u.link_key); } + else if (k->u.attach_count) + free (k->u.attach_count); if (aq) devicep->openacc.async.queue_callback_func (aq, gomp_unref_tgt_void, (void *) k->tgt); @@ -1372,6 +1544,18 @@ gomp_unmap_vars_internal (struct target_mem_desc *tgt, bool do_copyfrom, } size_t i; + + /* We must perform detachments before any copies back to the host. */ + for (i = 0; i < tgt->list_count; i++) + { + splay_tree_key k = tgt->list[i].key; + + if (k != NULL && tgt->list[i].do_detach) + gomp_detach_pointer (devicep, aq, k, tgt->list[i].key->host_start + + tgt->list[i].offset, + k->refcount == 1, NULL); + } + for (i = 0; i < tgt->list_count; i++) { splay_tree_key k = tgt->list[i].key; @@ -1545,6 +1729,7 @@ gomp_load_image_to_device (struct gomp_device_descr *devicep, unsigned version, k->tgt_offset = target_table[i].start; k->refcount = REFCOUNT_INFINITY; k->virtual_refcount = 0; + k->u.attach_count = NULL; array->left = NULL; array->right = NULL; splay_tree_insert (&devicep->mem_map, array); @@ -1577,6 +1762,7 @@ gomp_load_image_to_device (struct gomp_device_descr *devicep, unsigned version, k->tgt_offset = target_var->start; k->refcount = target_size & link_bit ? REFCOUNT_LINK : REFCOUNT_INFINITY; k->virtual_refcount = 0; + k->u.attach_count = NULL; array->left = NULL; array->right = NULL; splay_tree_insert (&devicep->mem_map, array); diff --git a/libgomp/testsuite/libgomp.oacc-c++/deep-copy-12.C b/libgomp/testsuite/libgomp.oacc-c++/deep-copy-12.C new file mode 100644 index 00000000000..771876afd2d --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c++/deep-copy-12.C @@ -0,0 +1,72 @@ +#include + +/* Test attach/detach with dereferences of reference to pointer to struct. */ + +typedef struct { + int *a; + int *b; + int *c; +} mystruct; + +int main(int argc, char* argv[]) +{ + const int N = 1024; + mystruct *m = (mystruct *) malloc (sizeof (*m)); + mystruct *&mref = m; + int i; + + mref->a = (int *) malloc (N * sizeof (int)); + m->b = (int *) malloc (N * sizeof (int)); + m->c = (int *) malloc (N * sizeof (int)); + + for (i = 0; i < N; i++) + { + mref->a[i] = 0; + m->b[i] = 0; + m->c[i] = 0; + } + +#pragma acc enter data copyin(m[0:1]) + + for (int i = 0; i < 99; i++) + { + int j; +#pragma acc parallel loop copy(mref->a[0:N]) + for (j = 0; j < N; j++) + mref->a[j]++; +#pragma acc parallel loop copy(mref->b[0:N], m->c[5:N-10]) + for (j = 0; j < N; j++) + { + mref->b[j]++; + if (j > 5 && j < N - 5) + m->c[j]++; + } + } + +#pragma acc exit data copyout(m[0:1]) + + for (i = 0; i < N; i++) + { + if (m->a[i] != 99) + abort (); + if (m->b[i] != 99) + abort (); + if (i > 5 && i < N-5) + { + if (m->c[i] != 99) + abort (); + } + else + { + if (m->c[i] != 0) + abort (); + } + } + + free (m->a); + free (m->b); + free (m->c); + free (m); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-c++/deep-copy-13.C b/libgomp/testsuite/libgomp.oacc-c++/deep-copy-13.C new file mode 100644 index 00000000000..98cf450c61d --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c++/deep-copy-13.C @@ -0,0 +1,72 @@ +#include + +/* Test array slice with reference to pointer. */ + +typedef struct { + int *a; + int *b; + int *c; +} mystruct; + +int main(int argc, char* argv[]) +{ + const int N = 1024; + mystruct *m = (mystruct *) malloc (sizeof (*m)); + int i; + + m->a = (int *) malloc (N * sizeof (int)); + m->b = (int *) malloc (N * sizeof (int)); + m->c = (int *) malloc (N * sizeof (int)); + + for (i = 0; i < N; i++) + { + m->a[i] = 0; + m->b[i] = 0; + m->c[i] = 0; + } + +#pragma acc enter data copyin(m[0:1]) + + for (int i = 0; i < 99; i++) + { + int j; + int *&ptr = m->a; +#pragma acc parallel loop copy(ptr[0:N]) + for (j = 0; j < N; j++) + ptr[j]++; +#pragma acc parallel loop copy(m->b[0:N], m->c[5:N-10]) + for (j = 0; j < N; j++) + { + m->b[j]++; + if (j > 5 && j < N - 5) + m->c[j]++; + } + } + +#pragma acc exit data copyout(m[0:1]) + + for (i = 0; i < N; i++) + { + if (m->a[i] != 99) + abort (); + if (m->b[i] != 99) + abort (); + if (i > 5 && i < N-5) + { + if (m->c[i] != 99) + abort (); + } + else + { + if (m->c[i] != 0) + abort (); + } + } + + free (m->a); + free (m->b); + free (m->c); + free (m); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-1.c new file mode 100644 index 00000000000..d8d7067e452 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-1.c @@ -0,0 +1,24 @@ +#include +#include + +struct dc +{ + int a; + int *b; +}; + +int +main () +{ + int n = 100, i; + struct dc v = { .a = 3, .b = (int *) malloc (sizeof (int) * n) }; + +#pragma acc parallel loop copy(v.a, v.b[:n]) + for (i = 0; i < n; i++) + v.b[i] = v.a; + + for (i = 0; i < 10; i++) + assert (v.b[i] == v.a); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-10.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-10.c new file mode 100644 index 00000000000..37cde4ef059 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-10.c @@ -0,0 +1,53 @@ +#include + +/* Test asyncronous attach and detach operation. */ + +typedef struct { + int *a; + int *b; +} mystruct; + +int +main (int argc, char* argv[]) +{ + const int N = 1024; + mystruct m; + int i; + + m.a = (int *) malloc (N * sizeof (int)); + m.b = (int *) malloc (N * sizeof (int)); + + for (i = 0; i < N; i++) + { + m.a[i] = 0; + m.b[i] = 0; + } + +#pragma acc enter data copyin(m) + + for (int i = 0; i < 99; i++) + { + int j; +#pragma acc parallel loop copy(m.a[0:N]) async(i % 2) + for (j = 0; j < N; j++) + m.a[j]++; +#pragma acc parallel loop copy(m.b[0:N]) async((i + 1) % 2) + for (j = 0; j < N; j++) + m.b[j]++; + } + +#pragma acc exit data copyout(m) wait(0, 1) + + for (i = 0; i < N; i++) + { + if (m.a[i] != 99) + abort (); + if (m.b[i] != 99) + abort (); + } + + free (m.a); + free (m.b); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-11.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-11.c new file mode 100644 index 00000000000..ed8ddcc54fa --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-11.c @@ -0,0 +1,72 @@ +#include + +/* Test multiple struct dereferences on one directive, and slices starting at + non-zero. */ + +typedef struct { + int *a; + int *b; + int *c; +} mystruct; + +int main(int argc, char* argv[]) +{ + const int N = 1024; + mystruct *m = (mystruct *) malloc (sizeof (*m)); + int i; + + m->a = (int *) malloc (N * sizeof (int)); + m->b = (int *) malloc (N * sizeof (int)); + m->c = (int *) malloc (N * sizeof (int)); + + for (i = 0; i < N; i++) + { + m->a[i] = 0; + m->b[i] = 0; + m->c[i] = 0; + } + +#pragma acc enter data copyin(m[0:1]) + + for (int i = 0; i < 99; i++) + { + int j; +#pragma acc parallel loop copy(m->a[0:N]) + for (j = 0; j < N; j++) + m->a[j]++; +#pragma acc parallel loop copy(m->b[0:N], m->c[5:N-10]) + for (j = 0; j < N; j++) + { + m->b[j]++; + if (j > 5 && j < N - 5) + m->c[j]++; + } + } + +#pragma acc exit data copyout(m[0:1]) + + for (i = 0; i < N; i++) + { + if (m->a[i] != 99) + abort (); + if (m->b[i] != 99) + abort (); + if (i > 5 && i < N-5) + { + if (m->c[i] != 99) + abort (); + } + else + { + if (m->c[i] != 0) + abort (); + } + } + + free (m->a); + free (m->b); + free (m->c); + free (m); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-14.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-14.c new file mode 100644 index 00000000000..04c70aed7b5 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-14.c @@ -0,0 +1,63 @@ +#include +#include + +/* Test attach/detach operation with chained dereferences. */ + +typedef struct mystruct { + int *a; + struct mystruct *next; +} mystruct; + +int +main (int argc, char* argv[]) +{ + const int N = 1024; + mystruct *m = (mystruct *) malloc (sizeof (*m)); + int i; + + m->a = (int *) malloc (N * sizeof (int)); + m->next = (mystruct *) malloc (sizeof (*m)); + m->next->a = (int *) malloc (N * sizeof (int)); + m->next->next = NULL; + + for (i = 0; i < N; i++) + { + m->a[i] = 0; + m->next->a[i] = 0; + } + +#pragma acc enter data copyin(m[0:1]) + acc_copyin (m->next, sizeof (*m)); + + for (int i = 0; i < 99; i++) + { + int j; + acc_copyin (m->next->a, N * sizeof (int)); + acc_attach ((void **) &m->next); + /* This will attach only the innermost pointer, i.e. "a[0:N]". That's + why we have to attach the "m->next" pointer manually above. */ +#pragma acc parallel loop copy(m->next->a[0:N]) + for (j = 0; j < N; j++) + m->next->a[j]++; + acc_detach ((void **) &m->next); + acc_copyout (m->next->a, N * sizeof (int)); + } + + acc_copyout (m->next, sizeof (*m)); +#pragma acc exit data copyout(m[0:1]) + + for (i = 0; i < N; i++) + { + if (m->a[i] != 0) + abort (); + if (m->next->a[i] != 99) + abort (); + } + + free (m->next->a); + free (m->next); + free (m->a); + free (m); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-2.c new file mode 100644 index 00000000000..7e26e9aa8b9 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-2.c @@ -0,0 +1,29 @@ +#include +#include + +int +main(int argc, char* argv[]) +{ + struct foo { + int *a, *b, c, d, *e; + } s; + + s.a = (int *) malloc (16 * sizeof (int)); + s.b = (int *) malloc (16 * sizeof (int)); + s.e = (int *) malloc (16 * sizeof (int)); + + #pragma acc data copy(s) + { + #pragma acc data copy(s.a[0:10]) + { + #pragma acc parallel loop attach(s.a) + for (int i = 0; i < 10; i++) + s.a[i] = i; + } + } + + for (int i = 0; i < 10; i++) + assert (s.a[i] == i); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-3.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-3.c new file mode 100644 index 00000000000..cec764bd3e7 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-3.c @@ -0,0 +1,34 @@ +#include +#include +#include + +int +main () +{ + int n = 100, i; + int *a = (int *) malloc (sizeof (int) * n); + int *b; + + for (i = 0; i < n; i++) + a[i] = i+1; + +#pragma acc enter data copyin(a[:n]) create(b) + + b = a; + acc_attach ((void **)&b); + +#pragma acc parallel loop present (b[:n]) + for (i = 0; i < n; i++) + b[i] = i+1; + + acc_detach ((void **)&b); + +#pragma acc exit data copyout(a[:n], b) + + for (i = 0; i < 10; i++) + assert (a[i] == b[i]); + + free (a); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-4.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-4.c new file mode 100644 index 00000000000..8874ca0a504 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-4.c @@ -0,0 +1,87 @@ +#include +#include + +#define LIST_LENGTH 10 + +struct node +{ + struct node *next; + int val; +}; + +int +sum_nodes (struct node *head) +{ + int i = 0, sum = 0; + +#pragma acc parallel reduction(+:sum) present(head[:1]) + { + for (; head != NULL; head = head->next) + sum += head->val; + } + + return sum; +} + +void +insert (struct node *head, int val) +{ + struct node *n = (struct node *) malloc (sizeof (struct node)); + + if (head->next) + { +#pragma acc exit data detach(head->next) + } + + n->val = val; + n->next = head->next; + head->next = n; + +#pragma acc enter data copyin(n[:1]) +#pragma acc enter data attach(head->next) + if (n->next) + { +#pragma acc enter data attach(n->next) + } +} + +void +destroy (struct node *head) +{ + while (head->next != NULL) + { +#pragma acc exit data detach(head->next) + struct node * n = head->next; + head->next = n->next; + if (n->next) + { +#pragma acc exit data detach(n->next) + } +#pragma acc exit data delete (n[:1]) + if (head->next) + { +#pragma acc enter data attach(head->next) + } + free (n); + } +} + +int +main () +{ + struct node list = { .next = NULL, .val = 0 }; + int i; + +#pragma acc enter data copyin(list) + + for (i = 0; i < LIST_LENGTH; i++) + insert (&list, i + 1); + + assert (sum_nodes (&list) == (LIST_LENGTH * LIST_LENGTH + LIST_LENGTH) / 2); + + destroy (&list); + +#pragma acc exit data delete(list) + + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-5.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-5.c new file mode 100644 index 00000000000..89cafbb62ab --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-5.c @@ -0,0 +1,81 @@ +#include +#include +#include + +struct node +{ + struct node *next; + int val; +}; + +int +sum_nodes (struct node *head) +{ + int i = 0, sum = 0; + +#pragma acc parallel reduction(+:sum) present(head[:1]) + { + for (; head != NULL; head = head->next) + sum += head->val; + } + + return sum; +} + +void +insert (struct node *head, int val) +{ + struct node *n = (struct node *) malloc (sizeof (struct node)); + + if (head->next) + acc_detach ((void **) &head->next); + + n->val = val; + n->next = head->next; + head->next = n; + + acc_copyin (n, sizeof (struct node)); + acc_attach((void **) &head->next); + + if (n->next) + acc_attach ((void **) &n->next); +} + +void +destroy (struct node *head) +{ + while (head->next != NULL) + { + acc_detach ((void **) &head->next); + struct node * n = head->next; + head->next = n->next; + if (n->next) + acc_detach ((void **) &n->next); + + acc_delete (n, sizeof (struct node)); + if (head->next) + acc_attach((void **) &head->next); + + free (n); + } +} + +int +main () +{ + struct node list = { .next = NULL, .val = 0 }; + int i; + + acc_copyin (&list, sizeof (struct node)); + + for (i = 0; i < 10; i++) + insert (&list, 2); + + assert (sum_nodes (&list) == 10 * 2); + + destroy (&list); + + acc_delete (&list, sizeof (struct node)); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-6.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-6.c new file mode 100644 index 00000000000..81c1c5e1093 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-6.c @@ -0,0 +1,59 @@ +/* { dg-do run { target { ! openacc_host_selected } } } */ + +#include +#include +#include + +struct dc +{ + int a; + int **b; +}; + +int +main () +{ + int n = 100, i, j, k; + struct dc v = { .a = 3 }; + + v.b = (int **) malloc (sizeof (int *) * n); + for (i = 0; i < n; i++) + v.b[i] = (int *) malloc (sizeof (int) * n); + + for (k = 0; k < 16; k++) + { +#pragma acc data copy(v) + { +#pragma acc data copy(v.b[:n]) + { + for (i = 0; i < n; i++) + { + acc_copyin (v.b[i], sizeof (int) * n); + acc_attach ((void **) &v.b[i]); + } + +#pragma acc parallel loop + for (i = 0; i < n; i++) + for (j = 0; j < n; j++) + v.b[i][j] = v.a + i + j; + + for (i = 0; i < n; i++) + { + acc_detach ((void **) &v.b[i]); + acc_copyout (v.b[i], sizeof (int) * n); + } + } + } + + for (i = 0; i < n; i++) + for (j = 0; j < n; j++) + assert (v.b[i][j] == v.a + i + j); + + assert (!acc_is_present (&v, sizeof (v))); + assert (!acc_is_present (v.b, sizeof (int *) * n)); + for (i = 0; i < n; i++) + assert (!acc_is_present (v.b[i], sizeof (int) * n)); + } + + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-7.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-7.c new file mode 100644 index 00000000000..a59047af520 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-7.c @@ -0,0 +1,45 @@ +/* { dg-do run { target { ! openacc_host_selected } } } */ + +#include +#include +#include + +struct dc +{ + int a; + int *b; +}; + +int +main () +{ + int n = 100, i, j, k; + struct dc v = { .a = 3 }; + + v.b = (int *) malloc (sizeof (int) * n); + + for (k = 0; k < 16; k++) + { + /* Here, we do not explicitly copy the enclosing structure, but work + with fields directly. Make sure attachment counters and reference + counters work properly in that case. */ +#pragma acc enter data copyin(v.a, v.b[0:n]) +#pragma acc enter data pcopyin(v.b[0:n]) +#pragma acc enter data pcopyin(v.b[0:n]) + +#pragma acc parallel loop present(v.a, v.b) + for (i = 0; i < n; i++) + v.b[i] = v.a + i; + +#pragma acc exit data copyout(v.b[:n]) finalize +#pragma acc exit data delete(v.a) + + for (i = 0; i < n; i++) + assert (v.b[i] == v.a + i); + + assert (!acc_is_present (&v, sizeof (v))); + assert (!acc_is_present (v.b, sizeof (int *) * n)); + } + + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-8.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-8.c new file mode 100644 index 00000000000..0ca5990b377 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-8.c @@ -0,0 +1,54 @@ +/* { dg-do run { target { ! openacc_host_selected } } } */ + +#include +#include +#include + +struct dc +{ + int a; + int *b; + int *c; + int *d; +}; + +int +main () +{ + int n = 100, i, j, k; + struct dc v = { .a = 3 }; + + v.b = (int *) malloc (sizeof (int) * n); + v.c = (int *) malloc (sizeof (int) * n); + v.d = (int *) malloc (sizeof (int) * n); + +#pragma acc enter data copyin(v) + + for (k = 0; k < 16; k++) + { +#pragma acc enter data copyin(v.a, v.b[:n], v.c[:n], v.d[:n]) + +#pragma acc parallel loop + for (i = 0; i < n; i++) + v.b[i] = v.a + i; + +#pragma acc exit data copyout(v.b[:n]) +#pragma acc exit data copyout(v.c[:n]) +#pragma acc exit data copyout(v.d[:n]) +#pragma acc exit data copyout(v.a) + + for (i = 0; i < n; i++) + assert (v.b[i] == v.a + i); + + assert (acc_is_present (&v, sizeof (v))); + assert (!acc_is_present (v.b, sizeof (int *) * n)); + assert (!acc_is_present (v.c, sizeof (int *) * n)); + assert (!acc_is_present (v.d, sizeof (int *) * n)); + } + +#pragma acc exit data copyout(v) + + assert (!acc_is_present (&v, sizeof (v))); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-9.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-9.c new file mode 100644 index 00000000000..28535c9e281 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-9.c @@ -0,0 +1,53 @@ +#include + +typedef struct { + int *a; + int *b; +} mystruct; + +int +main (int argc, char* argv[]) +{ + const int N = 1024; + mystruct *m = (mystruct *) malloc (sizeof (*m)); + int i; + + m->a = (int *) malloc (N * sizeof (int)); + m->b = (int *) malloc (N * sizeof (int)); + + for (i = 0; i < N; i++) + { + m->a[i] = 0; + m->b[i] = 0; + } + +#pragma acc enter data copyin(m[0:1]) + + for (int i = 0; i < 99; i++) + { + int j; + int *ptr = m->a; +#pragma acc parallel loop copy(m->a[0:N]) + for (j = 0; j < N; j++) + m->a[j]++; +#pragma acc parallel loop copy(m->b[0:N]) + for (j = 0; j < N; j++) + m->b[j]++; + } + +#pragma acc exit data copyout(m[0:1]) + + for (i = 0; i < N; i++) + { + if (m->a[i] != 99) + abort (); + if (m->b[i] != 99) + abort (); + } + + free (m->a); + free (m->b); + free (m); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-1.f90 new file mode 100644 index 00000000000..c4cea11b571 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-1.f90 @@ -0,0 +1,35 @@ +! { dg-do run } + +! Test of attach/detach with "acc data". + +program dtype + implicit none + integer, parameter :: n = 512 + type mytype + integer, allocatable :: a(:) + end type mytype + integer i + + type(mytype) :: var + + allocate(var%a(1:n)) + +!$acc data copy(var) +!$acc data copy(var%a) + +!$acc parallel loop + do i = 1,n + var%a(i) = i + end do +!$acc end parallel loop + +!$acc end data +!$acc end data + + do i = 1,n + if (i .ne. var%a(i)) stop 1 + end do + + deallocate(var%a) + +end program dtype diff --git a/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-2.f90 b/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-2.f90 new file mode 100644 index 00000000000..35936617b87 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-2.f90 @@ -0,0 +1,33 @@ +! { dg-do run } + +! Test of attach/detach with "acc data", two clauses at once. + +program dtype + implicit none + integer, parameter :: n = 512 + type mytype + integer, allocatable :: a(:) + end type mytype + integer i + + type(mytype) :: var + + allocate(var%a(1:n)) + +!$acc data copy(var) copy(var%a) + +!$acc parallel loop + do i = 1,n + var%a(i) = i + end do +!$acc end parallel loop + +!$acc end data + + do i = 1,n + if (i .ne. var%a(i)) stop 1 + end do + + deallocate(var%a) + +end program dtype diff --git a/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-3.f90 b/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-3.f90 new file mode 100644 index 00000000000..667d944fecb --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-3.f90 @@ -0,0 +1,34 @@ +! { dg-do run } + +! Test of attach/detach with "acc parallel". + +program dtype + implicit none + integer, parameter :: n = 512 + type mytype + integer, allocatable :: a(:) + integer, allocatable :: b(:) + end type mytype + integer i + + type(mytype) :: var + + allocate(var%a(1:n)) + allocate(var%b(1:n)) + +!$acc parallel loop copy(var) copy(var%a(1:n)) copy(var%b(1:n)) + do i = 1,n + var%a(i) = i + var%b(i) = i + end do +!$acc end parallel loop + + do i = 1,n + if (i .ne. var%a(i)) stop 1 + if (i .ne. var%b(i)) stop 2 + end do + + deallocate(var%a) + deallocate(var%b) + +end program dtype diff --git a/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-4.f90 b/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-4.f90 new file mode 100644 index 00000000000..6949e120c9f --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-4.f90 @@ -0,0 +1,49 @@ +! { dg-do run } + +! Test of attach/detach with "acc enter/exit data". + +program dtype + implicit none + integer, parameter :: n = 512 + type mytype + integer, allocatable :: a(:) + integer, allocatable :: b(:) + end type mytype + integer, allocatable :: r(:) + integer i + + type(mytype) :: var + + allocate(var%a(1:n)) + allocate(var%b(1:n)) + allocate(r(1:n)) + +!$acc enter data copyin(var) + +!$acc enter data copyin(var%a, var%b, r) + +!$acc parallel loop + do i = 1,n + var%a(i) = i + var%b(i) = i * 2 + r(i) = i * 3 + end do +!$acc end parallel loop + +!$acc exit data copyout(var%a) +!$acc exit data copyout(var%b) +!$acc exit data copyout(r) + + do i = 1,n + if (i .ne. var%a(i)) stop 1 + if (i * 2 .ne. var%b(i)) stop 2 + if (i * 3 .ne. r(i)) stop 3 + end do + +!$acc exit data delete(var) + + deallocate(var%a) + deallocate(var%b) + deallocate(r) + +end program dtype diff --git a/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-5.f90 b/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-5.f90 new file mode 100644 index 00000000000..6843cf1d0fa --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-5.f90 @@ -0,0 +1,57 @@ +! { dg-do run } + +! Test of attach/detach, "enter data" inside "data", and subarray. + +program dtype + implicit none + integer, parameter :: n = 512 + type mytype + integer, allocatable :: a(:) + integer, allocatable :: b(:) + end type mytype + integer i + + type(mytype) :: var + + allocate(var%a(1:n)) + allocate(var%b(1:n)) + +!$acc data copy(var) + + do i = 1, n + var%a(i) = 0 + var%b(i) = 0 + end do + +!$acc enter data copyin(var%a(5:n - 5), var%b(5:n - 5)) + +!$acc parallel loop + do i = 5,n - 5 + var%a(i) = i + var%b(i) = i * 2 + end do +!$acc end parallel loop + +!$acc exit data copyout(var%a(5:n - 5), var%b(5:n - 5)) + +!$acc end data + + do i = 1,4 + if (var%a(i) .ne. 0) stop 1 + if (var%b(i) .ne. 0) stop 2 + end do + + do i = 5,n - 5 + if (i .ne. var%a(i)) stop 3 + if (i * 2 .ne. var%b(i)) stop 4 + end do + + do i = n - 4,n + if (var%a(i) .ne. 0) stop 5 + if (var%b(i) .ne. 0) stop 6 + end do + + deallocate(var%a) + deallocate(var%b) + +end program dtype diff --git a/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-6.f90 b/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-6.f90 new file mode 100644 index 00000000000..12910d0d655 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-6.f90 @@ -0,0 +1,61 @@ +! { dg-do run } + +! Test of attachment counters and finalize. + +program dtype + implicit none + integer, parameter :: n = 512 + type mytype + integer, allocatable :: a(:) + integer, allocatable :: b(:) + end type mytype + integer i + + type(mytype) :: var + + allocate(var%a(1:n)) + allocate(var%b(1:n)) + +!$acc data copy(var) + + do i = 1, n + var%a(i) = 0 + var%b(i) = 0 + end do + +!$acc enter data copyin(var%a(5:n - 5), var%b(5:n - 5)) + + do i = 1,20 + !$acc enter data attach(var%a) + end do + +!$acc parallel loop + do i = 5,n - 5 + var%a(i) = i + var%b(i) = i * 2 + end do +!$acc end parallel loop + +!$acc exit data copyout(var%a(5:n - 5), var%b(5:n - 5)) finalize + +!$acc end data + + do i = 1,4 + if (var%a(i) .ne. 0) stop 1 + if (var%b(i) .ne. 0) stop 2 + end do + + do i = 5,n - 5 + if (i .ne. var%a(i)) stop 3 + if (i * 2 .ne. var%b(i)) stop 4 + end do + + do i = n - 4,n + if (var%a(i) .ne. 0) stop 5 + if (var%b(i) .ne. 0) stop 6 + end do + + deallocate(var%a) + deallocate(var%b) + +end program dtype diff --git a/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-7.f90 b/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-7.f90 new file mode 100644 index 00000000000..ab44f0a73b9 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-7.f90 @@ -0,0 +1,89 @@ +! { dg-do run } + +! Test of attach/detach with scalar elements and nested derived types. + +program dtype + implicit none + integer, parameter :: n = 512 + type subtype + integer :: g, h + integer, allocatable :: q(:) + end type subtype + type mytype + integer, allocatable :: a(:) + integer, allocatable :: c, d + integer, allocatable :: b(:) + integer :: f + type(subtype) :: s + end type mytype + integer i + + type(mytype) :: var + + allocate(var%a(1:n)) + allocate(var%b(1:n)) + allocate(var%c) + allocate(var%d) + allocate(var%s%q(1:n)) + + var%c = 16 + var%d = 20 + var%f = 7 + var%s%g = 21 + var%s%h = 38 + +!$acc enter data copyin(var) + + do i = 1, n + var%a(i) = 0 + var%b(i) = 0 + var%s%q(i) = 0 + end do + +!$acc data copy(var%a(5:n - 5), var%b(5:n - 5), var%c, var%d) & +!$acc & copy(var%s%q) + +!$acc parallel loop default(none) present(var) + do i = 5,n - 5 + var%a(i) = i + var%b(i) = i * 2 + var%s%q(i) = i * 3 + var%s%g = 100 + var%s%h = 101 + end do +!$acc end parallel loop + +!$acc end data + +!$acc exit data copyout(var) + + do i = 1,4 + if (var%a(i) .ne. 0) stop 1 + if (var%b(i) .ne. 0) stop 2 + if (var%s%q(i) .ne. 0) stop 3 + end do + + do i = 5,n - 5 + if (i .ne. var%a(i)) stop 4 + if (i * 2 .ne. var%b(i)) stop 5 + if (i * 3 .ne. var%s%q(i)) stop 6 + end do + + do i = n - 4,n + if (var%a(i) .ne. 0) stop 7 + if (var%b(i) .ne. 0) stop 8 + if (var%s%q(i) .ne. 0) stop 9 + end do + + if (var%c .ne. 16) stop 10 + if (var%d .ne. 20) stop 11 + if (var%s%g .ne. 100 .or. var%s%h .ne. 101) stop 12 + if (var%f .ne. 7) stop 13 + + deallocate(var%a) + deallocate(var%b) + deallocate(var%c) + deallocate(var%d) + deallocate(var%s%q) + +end program dtype diff --git a/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-8.f90 b/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-8.f90 new file mode 100644 index 00000000000..d142763ae59 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-8.f90 @@ -0,0 +1,41 @@ +! { dg-do run } + +! Test of explicit attach/detach clauses and attachment counters. There are no +! acc_attach/acc_detach API routines in Fortran. + +program dtype + use openacc + implicit none + integer, parameter :: n = 512 + type mytype + integer, allocatable :: a(:) + end type mytype + integer i + + type(mytype) :: var + + allocate(var%a(1:n)) + + call acc_copyin(var) + call acc_copyin(var%a) + + !$acc enter data attach(var%a) + +!$acc parallel loop attach(var%a) + do i = 1,n + var%a(i) = i + end do +!$acc end parallel loop + + !$acc exit data detach(var%a) + + call acc_copyout(var%a) + call acc_copyout(var) + + do i = 1,n + if (i .ne. var%a(i)) stop 1 + end do + + deallocate(var%a) + +end program dtype diff --git a/libgomp/testsuite/libgomp.oacc-fortran/derived-type-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/derived-type-1.f90 new file mode 100644 index 00000000000..eb7812d541e --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-fortran/derived-type-1.f90 @@ -0,0 +1,28 @@ +! Test derived types with subarrays + +! { dg-do run } + + implicit none + type dtype + integer :: a, b, c + end type dtype + integer, parameter :: n = 100 + integer i + type (dtype), dimension(n) :: d + + !$acc data copy(d(1:n)) + !$acc parallel loop + do i = 1, n + d(i)%a = i + d(i)%b = i-1 + d(i)%c = i+1 + end do + !$acc end data + + do i = 1, n + if (d(i)%a /= i) stop 1 + if (d(i)%b /= i-1) stop 2 + if (d(i)%c /= i+1) stop 3 + end do +end program + diff --git a/libgomp/testsuite/libgomp.oacc-fortran/derivedtype-1.f95 b/libgomp/testsuite/libgomp.oacc-fortran/derivedtype-1.f95 new file mode 100644 index 00000000000..75ce48ddca2 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-fortran/derivedtype-1.f95 @@ -0,0 +1,30 @@ +! { dg-do run } + +program main + implicit none + + type mytype + integer :: a, b, c + end type mytype + + type(mytype) :: myvar + integer :: i + + myvar%a = 0 + myvar%b = 0 + myvar%c = 0 + +!$acc enter data copyin(myvar) + +!$acc parallel present(myvar) + myvar%a = 1 + myvar%b = 2 + myvar%c = 3 +!$acc end parallel + +!$acc exit data copyout(myvar) + + if (myvar%a .ne. 1) stop 1 + if (myvar%b .ne. 2) stop 2 + if (myvar%c .ne. 3) stop 3 +end program main diff --git a/libgomp/testsuite/libgomp.oacc-fortran/derivedtype-2.f95 b/libgomp/testsuite/libgomp.oacc-fortran/derivedtype-2.f95 new file mode 100644 index 00000000000..3088b832957 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-fortran/derivedtype-2.f95 @@ -0,0 +1,41 @@ +! { dg-do run } + +program main + implicit none + + type tnest + integer :: ia, ib, ic + end type tnest + + type mytype + type(tnest) :: nest + integer :: a, b, c + end type mytype + + type(mytype) :: myvar + integer :: i + + myvar%a = 0 + myvar%b = 0 + myvar%c = 0 + myvar%nest%ia = 0 + myvar%nest%ib = 0 + myvar%nest%ic = 0 + +!$acc enter data copyin(myvar%nest) + +!$acc parallel present(myvar%nest) + myvar%nest%ia = 4 + myvar%nest%ib = 5 + myvar%nest%ic = 6 +!$acc end parallel + +!$acc exit data copyout(myvar%nest) + + if (myvar%a .ne. 0) stop 1 + if (myvar%b .ne. 0) stop 2 + if (myvar%c .ne. 0) stop 3 + if (myvar%nest%ia .ne. 4) stop 4 + if (myvar%nest%ib .ne. 5) stop 5 + if (myvar%nest%ic .ne. 6) stop 6 +end program main diff --git a/libgomp/testsuite/libgomp.oacc-fortran/multidim-slice.f95 b/libgomp/testsuite/libgomp.oacc-fortran/multidim-slice.f95 new file mode 100644 index 00000000000..a9b40eeab4c --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-fortran/multidim-slice.f95 @@ -0,0 +1,50 @@ +! { dg-do run } + +program main + implicit none + real, allocatable :: myarr(:,:,:,:,:) + integer i, j, k, l, m + + allocate(myarr(1:10,1:10,1:10,1:10,1:10)) + + do i=1,10 + do j=1,10 + do k=1,10 + do l=1,10 + do m=1,10 + myarr(m,l,k,j,i) = i+j+k+l+m + end do + end do + end do + end do + end do + + do i=1,10 + !$acc data copy(myarr(:,:,:,:,i)) + !$acc parallel loop collapse(4) present(myarr(:,:,:,:,i)) + do j=1,10 + do k=1,10 + do l=1,10 + do m=1,10 + myarr(m,l,k,j,i) = myarr(m,l,k,j,i) + 1 + end do + end do + end do + end do + !$acc end parallel loop + !$acc end data + end do + + do i=1,10 + do j=1,10 + do k=1,10 + do l=1,10 + do m=1,10 + if (myarr(m,l,k,j,i) .ne. i+j+k+l+m+1) stop 1 + end do + end do + end do + end do + end do + +end program main diff --git a/libgomp/testsuite/libgomp.oacc-fortran/update-2.f90 b/libgomp/testsuite/libgomp.oacc-fortran/update-2.f90 new file mode 100644 index 00000000000..c3c8a07868f --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-fortran/update-2.f90 @@ -0,0 +1,284 @@ +! Test ACC UPDATE with derived types. + +! { dg-do run } + +module dt + integer, parameter :: n = 10 + type inner + integer :: d(n) + end type inner + type mytype + integer(8) :: a, b, c(n) + type(inner) :: in + end type mytype +end module dt + +program derived_acc + use dt + + implicit none + integer i, res + type(mytype) :: var + + var%a = 0 + var%b = 1 + var%c(:) = 10 + var%in%d(:) = 100 + + var%c(:) = 10 + + !$acc enter data copyin(var) + + !$acc parallel loop present(var) + do i = 1, 1 + var%a = var%b + end do + !$acc end parallel loop + + !$acc update host(var%a) + + if (var%a /= var%b) stop 1 + + var%b = 100 + + !$acc update device(var%b) + + !$acc parallel loop present(var) + do i = 1, 1 + var%a = var%b + end do + !$acc end parallel loop + + !$acc update host(var%a) + + if (var%a /= var%b) stop 2 + + !$acc parallel loop present (var) + do i = 1, n + var%c(i) = i + end do + !$acc end parallel loop + + !$acc update host(var%c) + + var%a = -1 + + do i = 1, n + if (var%c(i) /= i) stop 3 + var%c(i) = var%a + end do + + !$acc update device(var%a) + !$acc update device(var%c) + + res = 0 + + !$acc parallel loop present(var) reduction(+:res) + do i = 1, n + if (var%c(i) /= var%a) res = res + 1 + end do + + if (res /= 0) stop 4 + + var%c(:) = 0 + + !$acc update device(var%c) + + !$acc parallel loop present(var) + do i = 5, 5 + var%c(i) = 1 + end do + !$acc end parallel loop + + !$acc update host(var%c(5)) + + do i = 1, n + if (i /= 5 .and. var%c(i) /= 0) stop 5 + if (i == 5 .and. var%c(i) /= 1) stop 6 + end do + + !$acc parallel loop present(var) + do i = 1, n + var%in%d = var%a + end do + !$acc end parallel loop + + !$acc update host(var%in%d) + + do i = 1, n + if (var%in%d(i) /= var%a) stop 7 + end do + + var%c(:) = 0 + + !$acc update device(var%c) + + var%c(:) = -1 + + !$acc parallel loop present(var) + do i = n/2, n + var%c(i) = i + end do + !$acc end parallel loop + + !$acc update host(var%c(n/2:n)) + + do i = 1,n + if (i < n/2 .and. var%c(i) /= -1) stop 8 + if (i >= n/2 .and. var%c(i) /= i) stop 9 + end do + + var%in%d(:) = 0 + !$acc update device(var%in%d) + + !$acc parallel loop present(var) + do i = 5, 5 + var%in%d(i) = 1 + end do + !$acc end parallel loop + + !$acc update host(var%in%d(5)) + + do i = 1, n + if (i /= 5 .and. var%in%d(i) /= 0) stop 10 + if (i == 5 .and. var%in%d(i) /= 1) stop 11 + end do + + !$acc exit data delete(var) + + call derived_acc_subroutine(var) +end program derived_acc + +subroutine derived_acc_subroutine(var) + use dt + + implicit none + integer i, res + type(mytype) :: var + + var%a = 0 + var%b = 1 + var%c(:) = 10 + var%in%d(:) = 100 + + var%c(:) = 10 + + !$acc enter data copyin(var) + + !$acc parallel loop present(var) + do i = 1, 1 + var%a = var%b + end do + !$acc end parallel loop + + !$acc update host(var%a) + + if (var%a /= var%b) stop 12 + + var%b = 100 + + !$acc update device(var%b) + + !$acc parallel loop present(var) + do i = 1, 1 + var%a = var%b + end do + !$acc end parallel loop + + !$acc update host(var%a) + + if (var%a /= var%b) stop 13 + + !$acc parallel loop present (var) + do i = 1, n + var%c(i) = i + end do + !$acc end parallel loop + + !$acc update host(var%c) + + var%a = -1 + + do i = 1, n + if (var%c(i) /= i) stop 14 + var%c(i) = var%a + end do + + !$acc update device(var%a) + !$acc update device(var%c) + + res = 0 + + !$acc parallel loop present(var) reduction(+:res) + do i = 1, n + if (var%c(i) /= var%a) res = res + 1 + end do + + if (res /= 0) stop 15 + + var%c(:) = 0 + + !$acc update device(var%c) + + !$acc parallel loop present(var) + do i = 5, 5 + var%c(i) = 1 + end do + !$acc end parallel loop + + !$acc update host(var%c(5)) + + do i = 1, n + if (i /= 5 .and. var%c(i) /= 0) stop 16 + if (i == 5 .and. var%c(i) /= 1) stop 17 + end do + + !$acc parallel loop present(var) + do i = 1, n + var%in%d = var%a + end do + !$acc end parallel loop + + !$acc update host(var%in%d) + + do i = 1, n + if (var%in%d(i) /= var%a) stop 18 + end do + + var%c(:) = 0 + + !$acc update device(var%c) + + var%c(:) = -1 + + !$acc parallel loop present(var) + do i = n/2, n + var%c(i) = i + end do + !$acc end parallel loop + + !$acc update host(var%c(n/2:n)) + + do i = 1,n + if (i < n/2 .and. var%c(i) /= -1) stop 19 + if (i >= n/2 .and. var%c(i) /= i) stop 20 + end do + + var%in%d(:) = 0 + !$acc update device(var%in%d) + + !$acc parallel loop present(var) + do i = 5, 5 + var%in%d(i) = 1 + end do + !$acc end parallel loop + + !$acc update host(var%in%d(5)) + + do i = 1, n + if (i /= 5 .and. var%in%d(i) /= 0) stop 21 + if (i == 5 .and. var%in%d(i) /= 1) stop 22 + end do + + !$acc exit data delete(var) +end subroutine derived_acc_subroutine