From patchwork Tue May 24 19:16:35 2016 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Cesar Philippidis X-Patchwork-Id: 625903 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@bilbo.ozlabs.org Received: from sourceware.org (server1.sourceware.org [209.132.180.131]) (using TLSv1.2 with cipher ECDHE-RSA-AES256-GCM-SHA384 (256/256 bits)) (No client certificate requested) by ozlabs.org (Postfix) with ESMTPS id 3rDlXL4h5Cz9sds for ; Wed, 25 May 2016 05:17:04 +1000 (AEST) Authentication-Results: ozlabs.org; dkim=pass (1024-bit key; unprotected) header.d=gcc.gnu.org header.i=@gcc.gnu.org header.b=xSQW3/mk; dkim-atps=neutral DomainKey-Signature: a=rsa-sha1; c=nofws; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender :subject:to:references:cc:from:message-id:date:mime-version :in-reply-to:content-type; q=dns; s=default; b=sLcBRFHNnjuRqQXFQ 5+qjHVggjV4spC1EDqYizzmwaBpMX8RcfpWAW5LpUrTqa8QASil7GAs1Jw+gd40n UxUF0aTi/pRmnW2U8FIHEsu+vOwQQuUOQRBH3dEd0MwuIJL+FwbpweOEEHnEh01j d7hMCh5EofqpzP6SWkG0StVUA0= 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 :subject:to:references:cc:from:message-id:date:mime-version :in-reply-to:content-type; s=default; bh=mdLbPPR4tEvh3BHDaLtC6QQ S/Og=; b=xSQW3/mkQrJlawLuGqy2Wg5LU4L+QKd/PkFx5TQE5SdkMqQNXjhrhsn F3AwFVJAjjvukp+RVfNIdTHldX9jX3Kul7Y3ax9YXLnZMJLLqd1FqJud0JgGn3YD nMrif9Uw98riIo2xzqciufX0EeyjS3JMJM0QBYJDR4DFK6eltFsU= Received: (qmail 61283 invoked by alias); 24 May 2016 19:16:51 -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 61259 invoked by uid 89); 24 May 2016 19:16:50 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-1.7 required=5.0 tests=AWL, BAYES_00, RCVD_IN_DNSWL_NONE, SPF_PASS autolearn=ham version=3.3.2 spammy=clique, UD:php, sum X-HELO: relay1.mentorg.com Received: from relay1.mentorg.com (HELO relay1.mentorg.com) (192.94.38.131) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with (AES256-GCM-SHA384 encrypted) ESMTPS; Tue, 24 May 2016 19:16:40 +0000 Received: from svr-orw-fem-05.mgc.mentorg.com ([147.34.97.43]) by relay1.mentorg.com with esmtp id 1b5Hoi-0006Fo-Eo from Cesar_Philippidis@mentor.com ; Tue, 24 May 2016 12:16:36 -0700 Received: from [127.0.0.1] (147.34.91.1) by svr-orw-fem-05.mgc.mentorg.com (147.34.97.43) with Microsoft SMTP Server id 14.3.224.2; Tue, 24 May 2016 12:16:36 -0700 Subject: Re: [patch, openacc] use firstprivate pointers for subarrays in c and c++ To: Jakub Jelinek References: <573244BE.5010708@codesourcery.com> <20160520094231.GF28550@tucnak.redhat.com> <5743BD19.4090708@codesourcery.com> <20160524060955.GB28550@tucnak.redhat.com> CC: "gcc-patches@gcc.gnu.org" From: Cesar Philippidis Message-ID: <5744A893.3040204@codesourcery.com> Date: Tue, 24 May 2016 12:16:35 -0700 User-Agent: Mozilla/5.0 (X11; Linux x86_64; rv:38.0) Gecko/20100101 Thunderbird/38.8.0 MIME-Version: 1.0 In-Reply-To: <20160524060955.GB28550@tucnak.redhat.com> On 05/23/2016 11:09 PM, Jakub Jelinek wrote: > On Mon, May 23, 2016 at 07:31:53PM -0700, Cesar Philippidis wrote: >> @@ -12559,7 +12560,7 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) >> t = OMP_CLAUSE_DECL (c); >> if (TREE_CODE (t) == TREE_LIST) >> { >> - if (handle_omp_array_sections (c, ort & C_ORT_OMP)) >> + if (handle_omp_array_sections (c, ort & (C_ORT_OMP | C_ORT_ACC))) >> { >> remove = true; >> break; > > You haven't touched the /c/ handle_omp_array_sections{,_1}. As I said, I believe > you can just drop the is_omp argument altogether (unlike C++), or, pass for > consistency ort itself there as well. But I bet the argument will be > unused. OK, I removed is_omp. I only had to guard one call to handle_omp_array_sections from c_finish_omp_clauses because OpenACC doesn't support array reductions. Is this OK for trunk? Cesar 2016-05-24 Cesar Philippidis gcc/c * c-parser.c (c_parser_oacc_declare): Add support for GOMP_MAP_FIRSTPRIVATE_POINTER. * c-typeck.c (handle_omp_array_sections_1): Remove is_omp argument. (handle_omp_array_sections): Likewise. (c_finish_omp_clauses): Add specific errors and warning messages for OpenACC. Use firsrtprivate pointers for OpenACC subarrays. Update calls to handle_omp_array_sections. gcc/cp/ * parser.c (cp_parser_oacc_declare): Add support for GOMP_MAP_FIRSTPRIVATE_POINTER. * semantics.c (handle_omp_array_sections_1): Replace bool is_omp argument with enum c_omp_region_type ort. Don't privatize OpenACC non-static members. (handle_omp_array_sections): Replace bool is_omp argument with enum c_omp_region_type ort. Update call to handle_omp_array_sections_1. (finish_omp_clauses): Add specific errors and warning messages for OpenACC. Use firsrtprivate pointers for OpenACC subarrays. Update call to handle_omp_array_sections. gcc/ * gimplify.c (omp_notice_variable): Use zero-length arrays for data pointers inside OACC_DATA regions. (gimplify_scan_omp_clauses): Prune firstprivate clause associated with OACC_DATA, OACC_ENTER_DATA and OACC_EXIT data regions. (gimplify_adjust_omp_clauses): Fix typo in comment. gcc/testsuite/ * c-c++-common/goacc/data-clause-duplicate-1.c: Adjust test. * c-c++-common/goacc/deviceptr-1.c: Likewise. * c-c++-common/goacc/kernels-alias-3.c: Likewise. * c-c++-common/goacc/kernels-alias-4.c: Likewise. * c-c++-common/goacc/kernels-alias-5.c: Likewise. * c-c++-common/goacc/kernels-alias-8.c: Likewise. * c-c++-common/goacc/kernels-alias-ipa-pta-3.c: Likewise. * c-c++-common/goacc/pcopy.c: Likewise. * c-c++-common/goacc/pcopyin.c: Likewise. * c-c++-common/goacc/pcopyout.c: Likewise. * c-c++-common/goacc/pcreate.c: Likewise. * c-c++-common/goacc/pr70688.c: New test. * c-c++-common/goacc/present-1.c: Adjust test. * c-c++-common/goacc/reduction-5.c: Likewise. * g++.dg/goacc/data-1.C: New test. libgomp/ * oacc-mem.c (acc_malloc): Update handling of shared-memory targets. (acc_free): Likewise. (acc_memcpy_to_device): Likewise. (acc_memcpy_from_device): Likewise. (acc_deviceptr): Likewise. (acc_hostptr): Likewise. (acc_is_present): Likewise. (acc_map_data): Likewise. (acc_unmap_data): Likewise. (present_create_copy): Likewise. (delete_copyout): Likewise. (update_dev_host): Likewise. * testsuite/libgomp.oacc-c-c++-common/asyncwait-1.c: Remove xfail. * testsuite/libgomp.oacc-c-c++-common/data-2-lib.c: New test. * testsuite/libgomp.oacc-c-c++-common/data-2.c: Adjust test. * testsuite/libgomp.oacc-c-c++-common/data-3.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/enter_exit-lib.c: New test. * testsuite/libgomp.oacc-c-c++-common/lib-13.c: Adjust test so that it only runs on nvptx targets. * testsuite/libgomp.oacc-c-c++-common/lib-14.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/lib-15.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/lib-16.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/lib-17.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/lib-18.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/lib-20.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/lib-21.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/lib-22.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/lib-23.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/lib-24.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/lib-25.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/lib-28.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/lib-29.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/lib-30.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/lib-34.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/lib-42.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/lib-43.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/lib-44.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/lib-47.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/lib-48.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/lib-52.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/lib-53.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/lib-54.c: Likewise. diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c index 1bc5eed..1cf4fb4 100644 --- a/gcc/c/c-parser.c +++ b/gcc/c/c-parser.c @@ -13602,6 +13602,7 @@ c_parser_oacc_declare (c_parser *parser) switch (OMP_CLAUSE_MAP_KIND (t)) { + case GOMP_MAP_FIRSTPRIVATE_POINTER: case GOMP_MAP_FORCE_ALLOC: case GOMP_MAP_FORCE_TO: case GOMP_MAP_FORCE_DEVICEPTR: diff --git a/gcc/c/c-typeck.c b/gcc/c/c-typeck.c index 74bad2a..33858bb 100644 --- a/gcc/c/c-typeck.c +++ b/gcc/c/c-typeck.c @@ -11939,8 +11939,7 @@ c_finish_omp_cancellation_point (location_t loc, tree clauses) static tree handle_omp_array_sections_1 (tree c, tree t, vec &types, - bool &maybe_zero_len, unsigned int &first_non_one, - bool is_omp) + bool &maybe_zero_len, unsigned int &first_non_one) { tree ret, low_bound, length, type; if (TREE_CODE (t) != TREE_LIST) @@ -11949,7 +11948,6 @@ handle_omp_array_sections_1 (tree c, tree t, vec &types, return error_mark_node; ret = t; if (TREE_CODE (t) == COMPONENT_REF - && is_omp && (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_TO || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FROM)) @@ -11996,7 +11994,7 @@ handle_omp_array_sections_1 (tree c, tree t, vec &types, } ret = handle_omp_array_sections_1 (c, TREE_CHAIN (t), types, - maybe_zero_len, first_non_one, is_omp); + maybe_zero_len, first_non_one); if (ret == error_mark_node || ret == NULL_TREE) return ret; @@ -12227,14 +12225,13 @@ handle_omp_array_sections_1 (tree c, tree t, vec &types, /* Handle array sections for clause C. */ static bool -handle_omp_array_sections (tree c, bool is_omp) +handle_omp_array_sections (tree c) { bool maybe_zero_len = false; unsigned int first_non_one = 0; auto_vec types; tree first = handle_omp_array_sections_1 (c, OMP_CLAUSE_DECL (c), types, - maybe_zero_len, first_non_one, - is_omp); + maybe_zero_len, first_non_one); if (first == error_mark_node) return true; if (first == NULL_TREE) @@ -12427,27 +12424,24 @@ handle_omp_array_sections (tree c, bool is_omp) && TREE_CODE (TREE_TYPE (t)) == ARRAY_TYPE)) return false; gcc_assert (OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_FORCE_DEVICEPTR); - if (is_omp) - switch (OMP_CLAUSE_MAP_KIND (c)) - { - case GOMP_MAP_ALLOC: - case GOMP_MAP_TO: - case GOMP_MAP_FROM: - case GOMP_MAP_TOFROM: - case GOMP_MAP_ALWAYS_TO: - case GOMP_MAP_ALWAYS_FROM: - case GOMP_MAP_ALWAYS_TOFROM: - case GOMP_MAP_RELEASE: - case GOMP_MAP_DELETE: - OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION (c) = 1; - break; - default: - break; - } + switch (OMP_CLAUSE_MAP_KIND (c)) + { + case GOMP_MAP_ALLOC: + case GOMP_MAP_TO: + case GOMP_MAP_FROM: + case GOMP_MAP_TOFROM: + case GOMP_MAP_ALWAYS_TO: + case GOMP_MAP_ALWAYS_FROM: + case GOMP_MAP_ALWAYS_TOFROM: + case GOMP_MAP_RELEASE: + case GOMP_MAP_DELETE: + OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION (c) = 1; + break; + default: + break; + } tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP); - if (!is_omp) - OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_POINTER); - else if (TREE_CODE (t) == COMPONENT_REF) + if (TREE_CODE (t) == COMPONENT_REF) OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ALWAYS_POINTER); else OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_FIRSTPRIVATE_POINTER); @@ -12520,7 +12514,7 @@ tree c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) { bitmap_head generic_head, firstprivate_head, lastprivate_head; - bitmap_head aligned_head, map_head, map_field_head; + bitmap_head aligned_head, map_head, map_field_head, oacc_reduction_head; tree c, t, type, *pc; tree simdlen = NULL_TREE, safelen = NULL_TREE; bool branch_seen = false; @@ -12537,6 +12531,7 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) bitmap_initialize (&aligned_head, &bitmap_default_obstack); bitmap_initialize (&map_head, &bitmap_default_obstack); bitmap_initialize (&map_field_head, &bitmap_default_obstack); + bitmap_initialize (&oacc_reduction_head, &bitmap_default_obstack); for (pc = &clauses, c = clauses; c ; c = *pc) { @@ -12560,7 +12555,7 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) t = OMP_CLAUSE_DECL (c); if (TREE_CODE (t) == TREE_LIST) { - if (handle_omp_array_sections (c, ort & C_ORT_OMP)) + if ((ort & C_ORT_OMP) && handle_omp_array_sections (c)) { remove = true; break; @@ -12874,6 +12869,17 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) omp_clause_code_name[OMP_CLAUSE_CODE (c)]); remove = true; } + else if (ort == C_ORT_ACC + && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_REDUCTION) + { + if (bitmap_bit_p (&oacc_reduction_head, DECL_UID (t))) + { + error ("%qD appears more than once in reduction clauses", t); + remove = true; + } + else + bitmap_set_bit (&oacc_reduction_head, DECL_UID (t)); + } else if (bitmap_bit_p (&generic_head, DECL_UID (t)) || bitmap_bit_p (&firstprivate_head, DECL_UID (t)) || bitmap_bit_p (&lastprivate_head, DECL_UID (t))) @@ -12885,7 +12891,10 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_PRIVATE && bitmap_bit_p (&map_head, DECL_UID (t))) { - error ("%qD appears both in data and map clauses", t); + if (ort == C_ORT_ACC) + error ("%qD appears more than once in data clauses", t); + else + error ("%qD appears both in data and map clauses", t); remove = true; } else @@ -12911,7 +12920,10 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) } else if (bitmap_bit_p (&map_head, DECL_UID (t))) { - error ("%qD appears both in data and map clauses", t); + if (ort == C_ORT_ACC) + error ("%qD appears more than once in data clauses", t); + else + error ("%qD appears both in data and map clauses", t); remove = true; } else @@ -13004,7 +13016,7 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) } if (TREE_CODE (t) == TREE_LIST) { - if (handle_omp_array_sections (c, ort & C_ORT_OMP)) + if (handle_omp_array_sections (c)) remove = true; break; } @@ -13027,7 +13039,7 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) t = OMP_CLAUSE_DECL (c); if (TREE_CODE (t) == TREE_LIST) { - if (handle_omp_array_sections (c, ort & C_ORT_OMP)) + if (handle_omp_array_sections (c)) remove = true; else { @@ -13054,6 +13066,9 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP) error ("%qD appears more than once in motion" " clauses", t); + else if (ort == C_ORT_ACC) + error ("%qD appears more than once in data" + " clauses", t); else error ("%qD appears more than once in map" " clauses", t); @@ -13155,7 +13170,10 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) } else if (bitmap_bit_p (&map_head, DECL_UID (t))) { - error ("%qD appears both in data and map clauses", t); + if (ort == C_ORT_ACC) + error ("%qD appears more than once in data clauses", t); + else + error ("%qD appears both in data and map clauses", t); remove = true; } else @@ -13165,6 +13183,8 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) { if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP) error ("%qD appears more than once in motion clauses", t); + else if (ort == C_ORT_ACC) + error ("%qD appears more than once in data clauses", t); else error ("%qD appears more than once in map clauses", t); remove = true; @@ -13172,7 +13192,10 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) else if (bitmap_bit_p (&generic_head, DECL_UID (t)) || bitmap_bit_p (&firstprivate_head, DECL_UID (t))) { - error ("%qD appears both in data and map clauses", t); + if (ort == C_ORT_ACC) + error ("%qD appears more than once in data clauses", t); + else + error ("%qD appears both in data and map clauses", t); remove = true; } else diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c index a68a510..d21230f 100644 --- a/gcc/cp/parser.c +++ b/gcc/cp/parser.c @@ -35214,6 +35214,7 @@ cp_parser_oacc_declare (cp_parser *parser, cp_token *pragma_tok) gcc_assert (OMP_CLAUSE_CODE (t) == OMP_CLAUSE_MAP); switch (OMP_CLAUSE_MAP_KIND (t)) { + case GOMP_MAP_FIRSTPRIVATE_POINTER: case GOMP_MAP_FORCE_ALLOC: case GOMP_MAP_FORCE_TO: case GOMP_MAP_FORCE_DEVICEPTR: diff --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c index 06dee5a..ad65a1b 100644 --- a/gcc/cp/semantics.c +++ b/gcc/cp/semantics.c @@ -4472,7 +4472,7 @@ omp_privatize_field (tree t, bool shared) static tree handle_omp_array_sections_1 (tree c, tree t, vec &types, bool &maybe_zero_len, unsigned int &first_non_one, - bool is_omp) + enum c_omp_region_type ort) { tree ret, low_bound, length, type; if (TREE_CODE (t) != TREE_LIST) @@ -4484,7 +4484,7 @@ handle_omp_array_sections_1 (tree c, tree t, vec &types, t = TREE_OPERAND (t, 0); ret = t; if (TREE_CODE (t) == COMPONENT_REF - && is_omp + && 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) @@ -4545,11 +4545,12 @@ handle_omp_array_sections_1 (tree c, tree t, vec &types, return ret; } - if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_REDUCTION + if (ort == C_ORT_OMP + && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_REDUCTION && TREE_CODE (TREE_CHAIN (t)) == FIELD_DECL) TREE_CHAIN (t) = omp_privatize_field (TREE_CHAIN (t), false); ret = handle_omp_array_sections_1 (c, TREE_CHAIN (t), types, - maybe_zero_len, first_non_one, is_omp); + maybe_zero_len, first_non_one, ort); if (ret == error_mark_node || ret == NULL_TREE) return ret; @@ -4792,14 +4793,14 @@ handle_omp_array_sections_1 (tree c, tree t, vec &types, /* Handle array sections for clause C. */ static bool -handle_omp_array_sections (tree c, bool is_omp) +handle_omp_array_sections (tree c, enum c_omp_region_type ort) { bool maybe_zero_len = false; unsigned int first_non_one = 0; auto_vec types; tree first = handle_omp_array_sections_1 (c, OMP_CLAUSE_DECL (c), types, maybe_zero_len, first_non_one, - is_omp); + ort); if (first == error_mark_node) return true; if (first == NULL_TREE) @@ -4988,7 +4989,7 @@ handle_omp_array_sections (tree c, bool is_omp) || (TREE_CODE (t) == COMPONENT_REF && TREE_CODE (TREE_TYPE (t)) == ARRAY_TYPE)) return false; - if (is_omp) + if (ort == C_ORT_OMP || ort == C_ORT_ACC) switch (OMP_CLAUSE_MAP_KIND (c)) { case GOMP_MAP_ALLOC: @@ -5007,7 +5008,7 @@ handle_omp_array_sections (tree c, bool is_omp) } tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP); - if (!is_omp) + 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); @@ -5774,7 +5775,7 @@ tree finish_omp_clauses (tree clauses, enum c_omp_region_type ort) { bitmap_head generic_head, firstprivate_head, lastprivate_head; - bitmap_head aligned_head, map_head, map_field_head; + bitmap_head aligned_head, map_head, map_field_head, oacc_reduction_head; tree c, t, *pc; tree safelen = NULL_TREE; bool branch_seen = false; @@ -5788,6 +5789,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) bitmap_initialize (&aligned_head, &bitmap_default_obstack); bitmap_initialize (&map_head, &bitmap_default_obstack); bitmap_initialize (&map_field_head, &bitmap_default_obstack); + bitmap_initialize (&oacc_reduction_head, &bitmap_default_obstack); for (pc = &clauses, c = clauses; c ; c = *pc) { @@ -5807,8 +5809,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) t = OMP_CLAUSE_DECL (c); if (TREE_CODE (t) == TREE_LIST) { - if (handle_omp_array_sections (c, ((ort & C_ORT_OMP_DECLARE_SIMD) - == C_ORT_OMP))) + if (handle_omp_array_sections (c, ort)) { remove = true; break; @@ -6018,6 +6019,17 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) omp_clause_code_name[OMP_CLAUSE_CODE (c)]); remove = true; } + else if (ort == C_ORT_ACC + && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_REDUCTION) + { + if (bitmap_bit_p (&oacc_reduction_head, DECL_UID (t))) + { + error ("%qD appears more than once in reduction clauses", t); + remove = true; + } + else + bitmap_set_bit (&oacc_reduction_head, DECL_UID (t)); + } else if (bitmap_bit_p (&generic_head, DECL_UID (t)) || bitmap_bit_p (&firstprivate_head, DECL_UID (t)) || bitmap_bit_p (&lastprivate_head, DECL_UID (t))) @@ -6028,7 +6040,10 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_PRIVATE && bitmap_bit_p (&map_head, DECL_UID (t))) { - error ("%qD appears both in data and map clauses", t); + if (ort == C_ORT_ACC) + error ("%qD appears more than once in data clauses", t); + else + error ("%qD appears both in data and map clauses", t); remove = true; } else @@ -6038,7 +6053,8 @@ 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)) + && t == OMP_CLAUSE_DECL (c) + && ort != C_ORT_ACC) { OMP_CLAUSE_DECL (c) = omp_privatize_field (t, (OMP_CLAUSE_CODE (c) @@ -6054,7 +6070,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 ("% allowed in OpenMP only in %" " clauses"); @@ -6081,7 +6097,10 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) } else if (bitmap_bit_p (&map_head, DECL_UID (t))) { - error ("%qD appears both in data and map clauses", t); + if (ort == C_ORT_ACC) + error ("%qD appears more than once in data clauses", t); + else + error ("%qD appears both in data and map clauses", t); remove = true; } else @@ -6529,8 +6548,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) } if (TREE_CODE (t) == TREE_LIST) { - if (handle_omp_array_sections (c, ((ort & C_ORT_OMP_DECLARE_SIMD) - == C_ORT_OMP))) + if (handle_omp_array_sections (c, ort)) remove = true; break; } @@ -6564,8 +6582,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) t = OMP_CLAUSE_DECL (c); if (TREE_CODE (t) == TREE_LIST) { - if (handle_omp_array_sections (c, ((ort & C_ORT_OMP_DECLARE_SIMD) - == C_ORT_OMP))) + if (handle_omp_array_sections (c, ort)) remove = true; else { @@ -6594,6 +6611,9 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP) error ("%qD appears more than once in motion" " clauses", t); + else if (ort == C_ORT_ACC) + error ("%qD appears more than once in data" + " clauses", t); else error ("%qD appears more than once in map" " clauses", t); @@ -6681,7 +6701,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) omp_clause_code_name[OMP_CLAUSE_CODE (c)]); remove = true; } - else if (t == current_class_ptr) + else if (ort != C_ORT_ACC && t == current_class_ptr) { error ("% allowed in OpenMP only in %" " clauses"); @@ -6730,7 +6750,10 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) } else if (bitmap_bit_p (&map_head, DECL_UID (t))) { - error ("%qD appears both in data and map clauses", t); + if (ort == C_ORT_ACC) + error ("%qD appears more than once in data clauses", t); + else + error ("%qD appears both in data and map clauses", t); remove = true; } else @@ -6740,6 +6763,8 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) { if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP) error ("%qD appears more than once in motion clauses", t); + if (ort == C_ORT_ACC) + error ("%qD appears more than once in data clauses", t); else error ("%qD appears more than once in map clauses", t); remove = true; @@ -6747,7 +6772,10 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) else if (bitmap_bit_p (&generic_head, DECL_UID (t)) || bitmap_bit_p (&firstprivate_head, DECL_UID (t))) { - error ("%qD appears both in data and map clauses", t); + if (ort == C_ORT_ACC) + error ("%qD appears more than once in data clauses", t); + else + error ("%qD appears both in data and map clauses", t); remove = true; } else diff --git a/gcc/gimplify.c b/gcc/gimplify.c index e702bc4..5cffdc1 100644 --- a/gcc/gimplify.c +++ b/gcc/gimplify.c @@ -6279,6 +6279,9 @@ omp_notice_variable (struct gimplify_omp_ctx *ctx, tree decl, bool in_code) error ("variable %qE declared in enclosing " "% region", DECL_NAME (decl)); nflags |= GOVD_MAP; + if (octx->region_type == ORT_ACC_DATA + && (n2->value & GOVD_MAP_0LEN_ARRAY)) + nflags |= GOVD_MAP_0LEN_ARRAY; goto found_outer; } } @@ -6854,9 +6857,14 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, { case OMP_TARGET: break; + case OACC_DATA: + if (TREE_CODE (TREE_TYPE (decl)) != ARRAY_TYPE) + break; 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) @@ -7310,6 +7318,10 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, omp_notice_variable (outer_ctx, t, true); } } + if (code == OACC_DATA + && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP + && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER) + flags |= GOVD_MAP_0LEN_ARRAY; omp_add_variable (ctx, decl, flags); if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_REDUCTION && OMP_CLAUSE_REDUCTION_PLACEHOLDER (c)) @@ -7568,6 +7580,10 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, gcc_unreachable (); } + if (code == OACC_DATA + && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP + && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER) + remove = true; if (remove) *list_p = OMP_CLAUSE_CHAIN (c); else @@ -8028,7 +8044,7 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p, break; } decl = OMP_CLAUSE_DECL (c); - /* Data clasues associated with acc parallel reductions must be + /* Data clauses associated with acc parallel reductions must be compatible with present_or_copy. Warn and adjust the clause if that is not the case. */ if (ctx->region_type == ORT_ACC_PARALLEL) diff --git a/gcc/testsuite/c-c++-common/goacc/data-clause-duplicate-1.c b/gcc/testsuite/c-c++-common/goacc/data-clause-duplicate-1.c index 7a1cf68..6245beb 100644 --- a/gcc/testsuite/c-c++-common/goacc/data-clause-duplicate-1.c +++ b/gcc/testsuite/c-c++-common/goacc/data-clause-duplicate-1.c @@ -2,12 +2,12 @@ void fun (void) { float *fp; -#pragma acc parallel copy(fp[0:2],fp[0:2]) /* { dg-error "'fp' appears more than once in map clauses" } */ +#pragma acc parallel copy(fp[0:2],fp[0:2]) /* { dg-error "'fp' appears more than once in data clauses" } */ ; -#pragma acc kernels present_or_copyin(fp[3]) present_or_copyout(fp[7:4]) /* { dg-error "'fp' appears more than once in map clauses" } */ +#pragma acc kernels present_or_copyin(fp[3]) present_or_copyout(fp[7:4]) /* { dg-error "'fp' appears more than once in data clauses" } */ ; -#pragma acc data create(fp[:10]) deviceptr(fp) /* { dg-error "'fp' appears more than once in map clauses" } */ +#pragma acc data create(fp[:10]) deviceptr(fp) /* { dg-error "'fp' appears more than once in data clauses" } */ ; -#pragma acc data create(fp) present(fp) /* { dg-error "'fp' appears more than once in map clauses" } */ +#pragma acc data create(fp) present(fp) /* { dg-error "'fp' appears more than once in data clauses" } */ ; } diff --git a/gcc/testsuite/c-c++-common/goacc/deviceptr-1.c b/gcc/testsuite/c-c++-common/goacc/deviceptr-1.c index 08ddb10..3aa0e8a 100644 --- a/gcc/testsuite/c-c++-common/goacc/deviceptr-1.c +++ b/gcc/testsuite/c-c++-common/goacc/deviceptr-1.c @@ -47,7 +47,7 @@ fun2 (void) /* { dg-error "'u' undeclared" "u undeclared" { target *-*-* } 46 } */ /* { dg-error "'fun2' is not a variable" "fun2 not a variable" { target *-*-* } 46 } */ /* { dg-error "'i' is not a pointer variable" "i not a pointer variable" { target *-*-* } 46 } */ - /* { dg-error "'fp' appears more than once in map clauses" "fp more than once" { target *-*-* } 46 } */ + /* { dg-error "'fp' appears more than once in data clauses" "fp more than once" { target *-*-* } 46 } */ ; } @@ -55,11 +55,11 @@ void fun3 (void) { float *fp; -#pragma acc data deviceptr(fp,fp) /* { dg-error "'fp' appears more than once in map clauses" } */ +#pragma acc data deviceptr(fp,fp) /* { dg-error "'fp' appears more than once in data clauses" } */ ; -#pragma acc parallel deviceptr(fp) deviceptr(fp) /* { dg-error "'fp' appears more than once in map clauses" } */ +#pragma acc parallel deviceptr(fp) deviceptr(fp) /* { dg-error "'fp' appears more than once in data clauses" } */ ; -#pragma acc kernels copy(fp) deviceptr(fp) /* { dg-error "'fp' appears more than once in map clauses" } */ +#pragma acc kernels copy(fp) deviceptr(fp) /* { dg-error "'fp' appears more than once in data clauses" } */ ; } diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-alias-3.c b/gcc/testsuite/c-c++-common/goacc/kernels-alias-3.c index 6989c1c..2934f12 100644 --- a/gcc/testsuite/c-c++-common/goacc/kernels-alias-3.c +++ b/gcc/testsuite/c-c++-common/goacc/kernels-alias-3.c @@ -17,5 +17,5 @@ foo (void) /* Only the omp_data_i related loads should be annotated with non-base 0 cliques. */ /* { dg-final { scan-tree-dump-times "clique 1 base 1" 2 "ealias" } } */ -/* { dg-final { scan-tree-dump-times "(?n)clique 1 base 0" 3 "ealias" } } */ +/* { dg-final { scan-tree-dump-times "(?n)clique 1 base 0" 2 "ealias" } } */ diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-alias-4.c b/gcc/testsuite/c-c++-common/goacc/kernels-alias-4.c index d41802c..f6ee5b5 100644 --- a/gcc/testsuite/c-c++-common/goacc/kernels-alias-4.c +++ b/gcc/testsuite/c-c++-common/goacc/kernels-alias-4.c @@ -19,5 +19,5 @@ foo (void) /* Only the omp_data_i related loads should be annotated with non-base 0 cliques. */ /* { dg-final { scan-tree-dump-times "clique 1 base 1" 2 "ealias" } } */ -/* { dg-final { scan-tree-dump-times "(?n)clique 1 base 0" 3 "ealias" } } */ +/* { dg-final { scan-tree-dump-times "(?n)clique 1 base 0" 2 "ealias" } } */ diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-alias-5.c b/gcc/testsuite/c-c++-common/goacc/kernels-alias-5.c index 6fefe183..74425fb 100644 --- a/gcc/testsuite/c-c++-common/goacc/kernels-alias-5.c +++ b/gcc/testsuite/c-c++-common/goacc/kernels-alias-5.c @@ -15,5 +15,5 @@ foo (int *a) /* Only the omp_data_i related loads should be annotated with cliques. */ /* { dg-final { scan-tree-dump-times "clique 1 base 1" 2 "ealias" } } */ -/* { dg-final { scan-tree-dump-times "(?n)clique 1 base 0" 4 "ealias" } } */ +/* { dg-final { scan-tree-dump-times "(?n)clique 1 base 0" 2 "ealias" } } */ diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-alias-8.c b/gcc/testsuite/c-c++-common/goacc/kernels-alias-8.c index 3b91acd..69200cc 100644 --- a/gcc/testsuite/c-c++-common/goacc/kernels-alias-8.c +++ b/gcc/testsuite/c-c++-common/goacc/kernels-alias-8.c @@ -7,7 +7,7 @@ extern void *acc_copyin (void *, size_t); void foo (int *a, size_t n) { - int *p = (int *)acc_copyin (&a, n); + int *p = (int *)acc_copyin (a, n); #pragma acc kernels deviceptr (p) pcopy(a[0:n]) { @@ -18,5 +18,5 @@ foo (int *a, size_t n) /* Only the omp_data_i related loads should be annotated with cliques. */ /* { dg-final { scan-tree-dump-times "clique 1 base 1" 2 "ealias" } } */ -/* { dg-final { scan-tree-dump-times "(?n)clique 1 base 0" 3 "ealias" } } */ +/* { dg-final { scan-tree-dump-times "(?n)clique 1 base 0" 2 "ealias" } } */ diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-alias-ipa-pta-3.c b/gcc/testsuite/c-c++-common/goacc/kernels-alias-ipa-pta-3.c index 1eb56eb..1ea0e73 100644 --- a/gcc/testsuite/c-c++-common/goacc/kernels-alias-ipa-pta-3.c +++ b/gcc/testsuite/c-c++-common/goacc/kernels-alias-ipa-pta-3.c @@ -31,6 +31,5 @@ foo (void) free (c); } -/* { dg-final { scan-tree-dump-times "(?n)= 0;$" 1 "optimized" } } */ -/* { dg-final { scan-tree-dump-times "(?n)= 1;$" 1 "optimized" } } */ -/* { dg-final { scan-tree-dump-times "(?n)= \\*a" 1 "optimized" } } */ +/* { dg-final { scan-tree-dump-times "(?n)= 0;$" 1 "optimized" { target c } } } */ +/* { dg-final { scan-tree-dump-times "(?n)= 1;$" 1 "optimized" { target c } } } */ diff --git a/gcc/testsuite/c-c++-common/goacc/pcopy.c b/gcc/testsuite/c-c++-common/goacc/pcopy.c index 02c4383..0e0aad5 100644 --- a/gcc/testsuite/c-c++-common/goacc/pcopy.c +++ b/gcc/testsuite/c-c++-common/goacc/pcopy.c @@ -7,4 +7,4 @@ f (char *cp) ; } -/* { dg-final { scan-tree-dump-times "#pragma acc parallel map\\(tofrom:\\*\\(cp \\+ 3\\) \\\[len: 5]\\) map\\(alloc:cp \\\[pointer assign, bias: 3]\\)" 1 "original" } } */ +/* { dg-final { scan-tree-dump-times "#pragma acc parallel map\\(tofrom:\\*\\(cp \\+ 3\\) \\\[len: 5]\\) map\\(firstprivate:cp \\\[pointer assign, bias: 3]\\)" 1 "original" } } */ diff --git a/gcc/testsuite/c-c++-common/goacc/pcopyin.c b/gcc/testsuite/c-c++-common/goacc/pcopyin.c index 10911fc..3085251 100644 --- a/gcc/testsuite/c-c++-common/goacc/pcopyin.c +++ b/gcc/testsuite/c-c++-common/goacc/pcopyin.c @@ -7,4 +7,4 @@ f (char *cp) ; } -/* { dg-final { scan-tree-dump-times "#pragma acc parallel map\\(to:\\*\\(cp \\+ 4\\) \\\[len: 6]\\) map\\(alloc:cp \\\[pointer assign, bias: 4]\\)" 1 "original" } } */ +/* { dg-final { scan-tree-dump-times "#pragma acc parallel map\\(to:\\*\\(cp \\+ 4\\) \\\[len: 6]\\) map\\(firstprivate:cp \\\[pointer assign, bias: 4]\\)" 1 "original" } } */ diff --git a/gcc/testsuite/c-c++-common/goacc/pcopyout.c b/gcc/testsuite/c-c++-common/goacc/pcopyout.c index 703ac2f..47c454c 100644 --- a/gcc/testsuite/c-c++-common/goacc/pcopyout.c +++ b/gcc/testsuite/c-c++-common/goacc/pcopyout.c @@ -7,4 +7,4 @@ f (char *cp) ; } -/* { dg-final { scan-tree-dump-times "#pragma acc parallel map\\(from:\\*\\(cp \\+ 5\\) \\\[len: 7]\\) map\\(alloc:cp \\\[pointer assign, bias: 5]\\)" 1 "original" } } */ +/* { dg-final { scan-tree-dump-times "#pragma acc parallel map\\(from:\\*\\(cp \\+ 5\\) \\\[len: 7]\\) map\\(firstprivate:cp \\\[pointer assign, bias: 5]\\)" 1 "original" } } */ diff --git a/gcc/testsuite/c-c++-common/goacc/pcreate.c b/gcc/testsuite/c-c++-common/goacc/pcreate.c index 00bf155..a403e5a 100644 --- a/gcc/testsuite/c-c++-common/goacc/pcreate.c +++ b/gcc/testsuite/c-c++-common/goacc/pcreate.c @@ -7,4 +7,4 @@ f (char *cp) ; } -/* { dg-final { scan-tree-dump-times "#pragma acc parallel map\\(alloc:\\*\\(cp \\+ 6\\) \\\[len: 8]\\) map\\(alloc:cp \\\[pointer assign, bias: 6]\\)" 1 "original" } } */ +/* { dg-final { scan-tree-dump-times "#pragma acc parallel map\\(alloc:\\*\\(cp \\+ 6\\) \\\[len: 8]\\) map\\(firstprivate:cp \\\[pointer assign, bias: 6]\\)" 1 "original" } } */ diff --git a/gcc/testsuite/c-c++-common/goacc/pr70688.c b/gcc/testsuite/c-c++-common/goacc/pr70688.c new file mode 100644 index 0000000..5a23665 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/pr70688.c @@ -0,0 +1,48 @@ +const int n = 100; + +int +private_reduction () +{ + int i, r; + + #pragma acc parallel + #pragma acc loop private (r) reduction (+:r) + for (i = 0; i < 100; i++) + r += 10; + + return r; +} + +int +parallel_reduction () +{ + int sum = 0; + int dummy = 0; + +#pragma acc data copy (dummy) + { +#pragma acc parallel num_gangs (10) copy (sum) reduction (+:sum) + { + int v = 5; + sum += 10 + v; + } + } + + return sum; +} + +int +main () +{ + int i, s = 0; + +#pragma acc parallel num_gangs (10) copy (s) reduction (+:s) + for (i = 0; i < n; i++) + s += i+1; + +#pragma acc parallel num_gangs (10) reduction (+:s) copy (s) + for (i = 0; i < n; i++) + s += i+1; + + return 0; +} diff --git a/gcc/testsuite/c-c++-common/goacc/present-1.c b/gcc/testsuite/c-c++-common/goacc/present-1.c index 7537948..51362b2 100644 --- a/gcc/testsuite/c-c++-common/goacc/present-1.c +++ b/gcc/testsuite/c-c++-common/goacc/present-1.c @@ -7,4 +7,4 @@ f (char *cp) ; } -/* { dg-final { scan-tree-dump-times "#pragma acc parallel map\\(force_present:\\*\\(cp \\+ 7\\) \\\[len: 9]\\) map\\(alloc:cp \\\[pointer assign, bias: 7]\\)" 1 "original" } } */ +/* { dg-final { scan-tree-dump-times "#pragma acc parallel map\\(force_present:\\*\\(cp \\+ 7\\) \\\[len: 9]\\) map\\(firstprivate:cp \\\[pointer assign, bias: 7]\\)" 1 "original" } } */ diff --git a/gcc/testsuite/c-c++-common/goacc/reduction-5.c b/gcc/testsuite/c-c++-common/goacc/reduction-5.c index 74daad3..dfdbab9 100644 --- a/gcc/testsuite/c-c++-common/goacc/reduction-5.c +++ b/gcc/testsuite/c-c++-common/goacc/reduction-5.c @@ -7,9 +7,9 @@ main(void) { int v1; -#pragma acc parallel reduction(+:v1) private(v1) /* { dg-error "appears more than once in data clauses" } */ +#pragma acc parallel reduction(+:v1) private(v1) /* { dg-error "invalid private reduction" } */ ; -#pragma acc parallel reduction(+:v1) firstprivate(v1) /* { dg-error "appears more than once in data clauses" } */ +#pragma acc parallel reduction(+:v1) firstprivate(v1) /* { dg-error "invalid private reduction" } */ ; return 0; diff --git a/gcc/testsuite/g++.dg/goacc/data-1.C b/gcc/testsuite/g++.dg/goacc/data-1.C new file mode 100644 index 0000000..54676dc --- /dev/null +++ b/gcc/testsuite/g++.dg/goacc/data-1.C @@ -0,0 +1,39 @@ +void +foo (int &a, int (&b)[100], int &n) +{ +#pragma acc enter data copyin (a, b) async wait +#pragma acc enter data create (b[20:30]) async wait +#pragma acc enter data (a) /* { dg-error "expected '#pragma acc' clause before '\\\(' token" } */ +#pragma acc enter data create (b(1:10)) /* { dg-error "expected '\\\)' before '\\\(' token" } */ +#pragma acc exit data delete (a) if (0) +#pragma acc exit data copyout (b) if (a) +#pragma acc exit data delete (b) +#pragma acc enter /* { dg-error "expected 'data' in" } */ +#pragma acc exit /* { dg-error "expected 'data' in" } */ +#pragma acc enter data /* { dg-error "has no data movement clause" } */ +#pragma acc exit data /* { dg-error "has no data movement clause" } */ +#pragma acc enter Data /* { dg-error "invalid pragma before" } */ +#pragma acc exit copyout (b) /* { dg-error "invalid pragma before" } */ +} + +template +void +foo (T &a, T (&b)[100], T &n) +{ +#pragma acc enter data copyin (a, b) async wait +#pragma acc enter data create (b[20:30]) async wait +#pragma acc enter data (a) /* { dg-error "expected '#pragma acc' clause before '\\\(' token" } */ +#pragma acc enter data create (b(1:10)) /* { dg-error "expected '\\\)' before '\\\(' token" } */ +#pragma acc exit data delete (a) if (0) +#pragma acc exit data copyout (b) if (a) +#pragma acc exit data delete (b) +#pragma acc enter /* { dg-error "expected 'data' in" } */ +#pragma acc exit /* { dg-error "expected 'data' in" } */ +#pragma acc enter data /* { dg-error "has no data movement clause" } */ +#pragma acc exit data /* { dg-error "has no data movement clause" } */ +#pragma acc enter Data /* { dg-error "invalid pragma before" } */ +#pragma acc exit copyout (b) /* { dg-error "invalid pragma before" } */ +} + +/* { dg-error "has no data movement clause" "" { target *-*-* } 6 } */ +/* { dg-error "has no data movement clause" "" { target *-*-* } 25 } */ diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c index ce1905c..665e208 100644 --- a/libgomp/oacc-mem.c +++ b/libgomp/oacc-mem.c @@ -32,6 +32,7 @@ #include "gomp-constants.h" #include "oacc-int.h" #include +#include #include /* Return block containing [H->S), or NULL if not contained. The device lock @@ -104,6 +105,9 @@ acc_malloc (size_t s) assert (thr->dev); + if (thr->dev->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM) + return malloc (s); + return thr->dev->alloc_func (thr->dev->target_id, s); } @@ -124,6 +128,9 @@ acc_free (void *d) struct gomp_device_descr *acc_dev = thr->dev; + if (acc_dev->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM) + return free (d); + gomp_mutex_lock (&acc_dev->lock); /* We don't have to call lazy open here, as the ptr value must have @@ -154,6 +161,12 @@ acc_memcpy_to_device (void *d, void *h, size_t s) assert (thr && thr->dev); + if (thr->dev->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM) + { + memmove (d, h, s); + return; + } + thr->dev->host2dev_func (thr->dev->target_id, d, h, s); } @@ -166,6 +179,12 @@ acc_memcpy_from_device (void *h, void *d, size_t s) assert (thr && thr->dev); + if (thr->dev->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM) + { + memmove (h, d, s); + return; + } + thr->dev->dev2host_func (thr->dev->target_id, h, d, s); } @@ -184,6 +203,9 @@ acc_deviceptr (void *h) struct goacc_thread *thr = goacc_thread (); struct gomp_device_descr *dev = thr->dev; + if (thr->dev->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM) + return h; + gomp_mutex_lock (&dev->lock); n = lookup_host (dev, h, 1); @@ -218,6 +240,9 @@ acc_hostptr (void *d) struct goacc_thread *thr = goacc_thread (); struct gomp_device_descr *acc_dev = thr->dev; + if (thr->dev->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM) + return d; + gomp_mutex_lock (&acc_dev->lock); n = lookup_dev (acc_dev->openacc.data_environ, d, 1); @@ -252,6 +277,9 @@ acc_is_present (void *h, size_t s) struct goacc_thread *thr = goacc_thread (); struct gomp_device_descr *acc_dev = thr->dev; + if (thr->dev->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM) + return h != NULL; + gomp_mutex_lock (&acc_dev->lock); n = lookup_host (acc_dev, h, s); @@ -271,7 +299,7 @@ acc_is_present (void *h, size_t s) void acc_map_data (void *h, void *d, size_t s) { - struct target_mem_desc *tgt; + struct target_mem_desc *tgt = NULL; size_t mapnum = 1; void *hostaddrs = h; void *devaddrs = d; @@ -287,9 +315,6 @@ acc_map_data (void *h, void *d, size_t s) { if (d != h) gomp_fatal ("cannot map data on shared-memory system"); - - tgt = gomp_map_vars (NULL, 0, NULL, NULL, NULL, NULL, true, - GOMP_MAP_VARS_OPENACC); } else { @@ -335,6 +360,10 @@ acc_unmap_data (void *h) /* No need to call lazy open, as the address must have been mapped. */ + /* This is a no-op on shared-memory targets. */ + if (acc_dev->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM) + return; + size_t host_size; gomp_mutex_lock (&acc_dev->lock); @@ -405,6 +434,9 @@ present_create_copy (unsigned f, void *h, size_t s) struct goacc_thread *thr = goacc_thread (); struct gomp_device_descr *acc_dev = thr->dev; + if (acc_dev->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM) + return h; + gomp_mutex_lock (&acc_dev->lock); n = lookup_host (acc_dev, h, s); @@ -496,6 +528,9 @@ delete_copyout (unsigned f, void *h, size_t s) struct goacc_thread *thr = goacc_thread (); struct gomp_device_descr *acc_dev = thr->dev; + if (acc_dev->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM) + return; + gomp_mutex_lock (&acc_dev->lock); n = lookup_host (acc_dev, h, s); @@ -553,6 +588,9 @@ update_dev_host (int is_dev, void *h, size_t s) struct goacc_thread *thr = goacc_thread (); struct gomp_device_descr *acc_dev = thr->dev; + if (acc_dev->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM) + return; + gomp_mutex_lock (&acc_dev->lock); n = lookup_host (acc_dev, h, s); diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/asyncwait-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/asyncwait-1.c index f3b490a..d478ce2 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/asyncwait-1.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/asyncwait-1.c @@ -1,6 +1,4 @@ /* { dg-do run { target openacc_nvidia_accel_selected } } */ -/* . - { dg-xfail-run-if "TODO" { *-*-* } } */ /* { dg-additional-options "-lcuda" } */ #include diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/data-2-lib.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/data-2-lib.c new file mode 100644 index 0000000..e1aa2c9 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/data-2-lib.c @@ -0,0 +1,185 @@ +/* This test is similar to data-2.c, but it uses acc_* library functions + to move data. */ + +/* { dg-do run } */ + +#include +#include +#include + +int +main (int argc, char **argv) +{ + int N = 128; //1024 * 1024; + float *a, *b, *c, *d, *e; + void *d_a, *d_b, *d_c, *d_d; + int i; + int nbytes; + + nbytes = N * sizeof (float); + + a = (float *) malloc (nbytes); + b = (float *) malloc (nbytes); + c = (float *) malloc (nbytes); + d = (float *) malloc (nbytes); + e = (float *) malloc (nbytes); + + for (i = 0; i < N; i++) + { + a[i] = 3.0; + b[i] = 0.0; + } + + d_a = acc_copyin (a, nbytes); + d_b = acc_copyin (b, nbytes); + acc_copyin (&N, sizeof (int)); + +#pragma acc parallel present (a[0:N], b[0:N], N) async wait +#pragma acc loop + for (i = 0; i < N; i++) + b[i] = a[i]; + + acc_wait_all (); + + acc_memcpy_from_device (a, d_a, nbytes); + acc_memcpy_from_device (b, d_b, nbytes); + + for (i = 0; i < N; i++) + { + assert (a[i] == 3.0); + assert (b[i] == 3.0); + } + + for (i = 0; i < N; i++) + { + a[i] = 2.0; + b[i] = 0.0; + } + + acc_update_device (a, nbytes); + acc_update_device (b, nbytes); + +#pragma acc parallel present (a[0:N], b[0:N], N) async (1) +#pragma acc loop + for (i = 0; i < N; i++) + b[i] = a[i]; + + acc_memcpy_from_device (a, d_a, nbytes); + acc_memcpy_from_device (b, d_b, nbytes); + + for (i = 0; i < N; i++) + { + assert (a[i] == 2.0); + assert (b[i] == 2.0); + } + + for (i = 0; i < N; i++) + { + a[i] = 3.0; + b[i] = 0.0; + c[i] = 0.0; + d[i] = 0.0; + } + + acc_update_device (a, nbytes); + acc_update_device (b, nbytes); + d_c = acc_copyin (c, nbytes); + d_d = acc_copyin (d, nbytes); + +#pragma acc parallel present (a[0:N], b[0:N], N) async (1) +#pragma acc loop + for (i = 0; i < N; i++) + b[i] = (a[i] * a[i] * a[i]) / a[i]; + +#pragma acc parallel present (a[0:N], c[0:N], N) async (2) +#pragma acc loop + for (i = 0; i < N; i++) + c[i] = (a[i] + a[i] + a[i] + a[i]) / a[i]; + +#pragma acc parallel present (a[0:N], d[0:N], N) async (3) +#pragma acc loop + for (i = 0; i < N; i++) + d[i] = ((a[i] * a[i] + a[i]) / a[i]) - a[i]; + + acc_wait_all (); + + acc_memcpy_from_device (a, d_a, nbytes); + acc_memcpy_from_device (b, d_b, nbytes); + acc_memcpy_from_device (c, d_c, nbytes); + acc_memcpy_from_device (d, d_d, nbytes); + + for (i = 0; i < N; i++) + { + if (a[i] != 3.0) + abort (); + + if (b[i] != 9.0) + abort (); + + if (c[i] != 4.0) + abort (); + + if (d[i] != 1.0) + abort (); + } + + for (i = 0; i < N; i++) + { + a[i] = 2.0; + b[i] = 0.0; + c[i] = 0.0; + d[i] = 0.0; + e[i] = 0.0; + } + + acc_update_device (a, nbytes); + acc_update_device (b, nbytes); + acc_update_device (c, nbytes); + acc_update_device (d, nbytes); + acc_copyin (e, nbytes); + +#pragma acc parallel present (a[0:N], b[0:N], N) async (1) + for (int ii = 0; ii < N; ii++) + b[ii] = (a[ii] * a[ii] * a[ii]) / a[ii]; + +#pragma acc parallel present (a[0:N], c[0:N], N) async (2) + for (int ii = 0; ii < N; ii++) + c[ii] = (a[ii] + a[ii] + a[ii] + a[ii]) / a[ii]; + +#pragma acc parallel present (a[0:N], d[0:N], N) async (3) + for (int ii = 0; ii < N; ii++) + d[ii] = ((a[ii] * a[ii] + a[ii]) / a[ii]) - a[ii]; + +#pragma acc parallel present (a[0:N], b[0:N], c[0:N], d[0:N], e[0:N], N) \ + async (4) + for (int ii = 0; ii < N; ii++) + e[ii] = a[ii] + b[ii] + c[ii] + d[ii]; + + acc_wait_all (); + acc_copyout (a, nbytes); + acc_copyout (b, nbytes); + acc_copyout (c, nbytes); + acc_copyout (d, nbytes); + acc_copyout (e, nbytes); + acc_delete (&N, sizeof (int)); + + for (i = 0; i < N; i++) + { + if (a[i] != 2.0) + abort (); + + if (b[i] != 4.0) + abort (); + + if (c[i] != 4.0) + abort (); + + if (d[i] != 1.0) + abort (); + + if (e[i] != 11.0) + abort (); + } + + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/data-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/data-2.c index f867a66..c1c0825 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/data-2.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/data-2.c @@ -1,3 +1,5 @@ +/* Test 'acc enter/exit data' regions. */ + /* { dg-do run } */ #include @@ -25,7 +27,7 @@ main (int argc, char **argv) } #pragma acc enter data copyin (a[0:N]) copyin (b[0:N]) copyin (N) async -#pragma acc parallel async wait +#pragma acc parallel present (a[0:N], b[0:N]) async wait #pragma acc loop for (i = 0; i < N; i++) b[i] = a[i]; @@ -49,7 +51,7 @@ main (int argc, char **argv) } #pragma acc enter data copyin (a[0:N]) copyin (b[0:N]) copyin (N) async (1) -#pragma acc parallel async (1) +#pragma acc parallel present (a[0:N], b[0:N]) async (1) #pragma acc loop for (i = 0; i < N; i++) b[i] = a[i]; @@ -76,17 +78,17 @@ main (int argc, char **argv) #pragma acc enter data copyin (a[0:N]) copyin (b[0:N]) copyin (c[0:N]) copyin (d[0:N]) copyin (N) async (1) -#pragma acc parallel async (1) wait (1) +#pragma acc parallel present (a[0:N], b[0:N]) async (1) wait (1) #pragma acc loop for (i = 0; i < N; i++) b[i] = (a[i] * a[i] * a[i]) / a[i]; -#pragma acc parallel async (2) wait (1) +#pragma acc parallel present (a[0:N], c[0:N]) async (2) wait (1) #pragma acc loop for (i = 0; i < N; i++) c[i] = (a[i] + a[i] + a[i] + a[i]) / a[i]; -#pragma acc parallel async (3) wait (1) +#pragma acc parallel present (a[0:N], d[0:N]) async (3) wait (1) #pragma acc loop for (i = 0; i < N; i++) d[i] = ((a[i] * a[i] + a[i]) / a[i]) - a[i]; @@ -120,26 +122,27 @@ main (int argc, char **argv) #pragma acc enter data copyin (a[0:N]) copyin (b[0:N]) copyin (c[0:N]) copyin (d[0:N]) copyin (e[0:N]) copyin (N) async (1) -#pragma acc parallel async (1) wait (1) +#pragma acc parallel present (a[0:N], b[0:N]) async (1) wait (1) for (int ii = 0; ii < N; ii++) b[ii] = (a[ii] * a[ii] * a[ii]) / a[ii]; -#pragma acc parallel async (2) wait (1) +#pragma acc parallel present (a[0:N], c[0:N]) async (2) wait (1) for (int ii = 0; ii < N; ii++) c[ii] = (a[ii] + a[ii] + a[ii] + a[ii]) / a[ii]; -#pragma acc parallel async (3) wait (1) +#pragma acc parallel present (a[0:N], d[0:N]) async (3) wait (1) for (int ii = 0; ii < N; ii++) d[ii] = ((a[ii] * a[ii] + a[ii]) / a[ii]) - a[ii]; -#pragma acc parallel wait (1) async (4) +#pragma acc parallel present (a[0:N], b[0:N], c[0:N], d[0:N], e[0:N]) \ + wait (1) async (4) for (int ii = 0; ii < N; ii++) e[ii] = a[ii] + b[ii] + c[ii] + d[ii]; -#pragma acc exit data copyout (a[0:N]) copyout (b[0:N]) copyout (c[0:N]) copyout (d[0:N]) copyout (e[0:N]) wait (1, 2, 3, 4) async (1) +#pragma acc exit data copyout (a[0:N]) copyout (b[0:N]) copyout (c[0:N]) \ + copyout (d[0:N]) copyout (e[0:N]) wait (1, 2, 3, 4) async (1) #pragma acc wait (1) - for (i = 0; i < N; i++) { if (a[i] != 2.0) diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/data-3.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/data-3.c index 747109f..0bf706a 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/data-3.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/data-3.c @@ -1,3 +1,5 @@ +/* Test 'acc enter/exit data' regions with 'acc update'. */ + /* { dg-do run } */ #include @@ -25,7 +27,7 @@ main (int argc, char **argv) } #pragma acc enter data copyin (a[0:N]) copyin (b[0:N]) copyin (N) async -#pragma acc parallel async wait +#pragma acc parallel present (a[0:N], b[0:N]) async wait #pragma acc loop for (i = 0; i < N; i++) b[i] = a[i]; @@ -49,7 +51,7 @@ main (int argc, char **argv) } #pragma acc update device (a[0:N], b[0:N]) async (1) -#pragma acc parallel async (1) +#pragma acc parallel present (a[0:N], b[0:N]) async (1) #pragma acc loop for (i = 0; i < N; i++) b[i] = a[i]; @@ -78,17 +80,17 @@ main (int argc, char **argv) #pragma acc update device (b[0:N]) async (2) #pragma acc enter data copyin (c[0:N], d[0:N]) async (3) -#pragma acc parallel async (1) wait (1,2) +#pragma acc parallel present (a[0:N], b[0:N]) async (1) wait (1,2) #pragma acc loop for (i = 0; i < N; i++) b[i] = (a[i] * a[i] * a[i]) / a[i]; -#pragma acc parallel async (2) wait (1,3) +#pragma acc parallel present (a[0:N], c[0:N]) async (2) wait (1,3) #pragma acc loop for (i = 0; i < N; i++) c[i] = (a[i] + a[i] + a[i] + a[i]) / a[i]; -#pragma acc parallel async (3) wait (1,3) +#pragma acc parallel present (a[0:N], d[0:N]) async (3) wait (1,3) #pragma acc loop for (i = 0; i < N; i++) d[i] = ((a[i] * a[i] + a[i]) / a[i]) - a[i]; @@ -123,27 +125,28 @@ main (int argc, char **argv) #pragma acc update device (a[0:N], b[0:N], c[0:N], d[0:N]) async (1) #pragma acc enter data copyin (e[0:N]) async (5) -#pragma acc parallel async (1) wait (1) +#pragma acc parallel present (a[0:N], b[0:N]) async (1) wait (1) for (int ii = 0; ii < N; ii++) b[ii] = (a[ii] * a[ii] * a[ii]) / a[ii]; -#pragma acc parallel async (2) wait (1) +#pragma acc parallel present (a[0:N], c[0:N]) async (2) wait (1) for (int ii = 0; ii < N; ii++) c[ii] = (a[ii] + a[ii] + a[ii] + a[ii]) / a[ii]; -#pragma acc parallel async (3) wait (1) +#pragma acc parallel present (a[0:N], d[0:N]) async (3) wait (1) for (int ii = 0; ii < N; ii++) d[ii] = ((a[ii] * a[ii] + a[ii]) / a[ii]) - a[ii]; -#pragma acc parallel wait (1,5) async (4) +#pragma acc parallel present (a[0:N], b[0:N], c[0:N], d[0:N], e[0:N]) \ + wait (1,5) async (4) for (int ii = 0; ii < N; ii++) e[ii] = a[ii] + b[ii] + c[ii] + d[ii]; -#pragma acc exit data copyout (a[0:N]) copyout (b[0:N]) copyout (c[0:N]) copyout (d[0:N]) copyout (e[0:N]) wait (1, 2, 3, 4) async (1) +#pragma acc exit data copyout (a[0:N]) copyout (b[0:N]) copyout (c[0:N]) \ + copyout (d[0:N]) copyout (e[0:N]) wait (1, 2, 3, 4) async (1) #pragma acc exit data delete (N) #pragma acc wait (1) - for (i = 0; i < N; i++) { if (a[i] != 2.0) @@ -162,5 +165,11 @@ main (int argc, char **argv) abort (); } + free (a); + free (b); + free (c); + free (d); + free (e); + return 0; } diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/enter_exit-lib.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/enter_exit-lib.c new file mode 100644 index 0000000..b5b37b2 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/enter_exit-lib.c @@ -0,0 +1,70 @@ +/* Verify enter/exit data interoperablilty between pragmas and + acc library calls. */ + +/* { dg-do run } */ + +#include +#include +#include + +int +main () +{ + int *p = (int *)malloc (sizeof (int)); + + /* Test 1: pragma input, library output. */ + +#pragma acc enter data copyin (p[0:1]) + +#pragma acc parallel present (p[0:1]) num_gangs (1) + { + p[0] = 1; + } + + acc_copyout (p, sizeof (int)); + + assert (p[0] == 1); + + /* Test 2: library input, pragma output. */ + + acc_copyin (p, sizeof (int)); + +#pragma acc parallel present (p[0:1]) num_gangs (1) + { + p[0] = 2; + } + +#pragma acc exit data copyout (p[0:1]) + + assert (p[0] == 2); + + /* Test 3: library input, library output. */ + + acc_copyin (p, sizeof (int)); + +#pragma acc parallel present (p[0:1]) num_gangs (1) + { + p[0] = 3; + } + + acc_copyout (p, sizeof (int)); + + assert (p[0] == 3); + + /* Test 4: pragma input, pragma output. */ + +#pragma acc enter data copyin (p[0:1]) + +#pragma acc parallel present (p[0:1]) num_gangs (1) + { + p[0] = 3; + } + +#pragma acc exit data copyout (p[0:1]) + + assert (p[0] == 3); + + free (p); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-13.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-13.c index 7098ef3..d665533 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-13.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-13.c @@ -1,4 +1,6 @@ -/* { dg-do run } */ +/* Check acc_is_present and acc_delete. */ + +/* { dg-do run { target openacc_nvidia_accel_selected } } */ #include #include diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-14.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-14.c index a9632f7..ee21257 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-14.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-14.c @@ -1,4 +1,6 @@ -/* { dg-do run } */ +/* Check acc_is_present. */ + +/* { dg-do run { target openacc_nvidia_accel_selected } } */ #include #include diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-15.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-15.c index 4f6a731..50c1701 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-15.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-15.c @@ -1,4 +1,6 @@ -/* { dg-do run } */ +/* Check acc_is_present and acc_copyout. */ + +/* { dg-do run { target openacc_nvidia_accel_selected } } */ #include #include diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-16.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-16.c index 28e4e5c..c81a78d 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-16.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-16.c @@ -1,4 +1,6 @@ -/* { dg-do run } */ +/* Test if duplicate data mappings with acc_copy_in. */ + +/* { dg-do run { target openacc_nvidia_accel_selected } } */ #include #include diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-17.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-17.c index 7d1767e..a3487e8 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-17.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-17.c @@ -1,4 +1,7 @@ -/* { dg-do run } */ +/* Check acc_copyout failure with acc_device_nvidia. */ + +/* { dg-do run { target openacc_nvidia_accel_selected } } */ + #include #include diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-18.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-18.c index 160b33c..b686cc9 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-18.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-18.c @@ -1,4 +1,6 @@ -/* { dg-do run } */ +/* Verify that acc_delete unregisters data mappings on the device. */ + +/* { dg-do run { target openacc_nvidia_accel_selected } } */ #include #include diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-20.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-20.c index 4f8e14c..25ceb3a 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-20.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-20.c @@ -1,4 +1,6 @@ -/* { dg-do run } */ +/* Exercise acc_copyin and acc_copyout on nvidia targets. */ + +/* { dg-do run { target openacc_nvidia_accel_selected } } */ #include #include diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-21.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-21.c index d908700..b170f81 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-21.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-21.c @@ -1,4 +1,6 @@ -/* { dg-do run } */ +/* Exercise acc_copyin and acc_copyout on nvidia targets. */ + +/* { dg-do run { target openacc_nvidia_accel_selected } } */ #include #include diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-22.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-22.c index a6c0197..65ff440 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-22.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-22.c @@ -1,4 +1,6 @@ -/* { dg-do run } */ +/* Exercise acc_copyin and acc_copyout on nvidia targets. */ + +/* { dg-do run { target openacc_nvidia_accel_selected } } */ #include #include diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-23.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-23.c index 2339dd6..fd4dc59 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-23.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-23.c @@ -1,4 +1,6 @@ -/* { dg-do run } */ +/* Exercise acc_copyin and acc_copyout on nvidia targets. */ + +/* { dg-do run { target openacc_nvidia_accel_selected } } */ #include #include diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-24.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-24.c index d7de8e3..09e2817 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-24.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-24.c @@ -1,4 +1,6 @@ -/* { dg-do run } */ +/* Exercise acc_create, acc_is_present and acc_delete. */ + +/* { dg-do run { target openacc_nvidia_accel_selected } } */ #include #include diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-25.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-25.c index bb709d3..5f00ccb 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-25.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-25.c @@ -1,4 +1,6 @@ -/* { dg-do run } */ +/* Exercise acc_create and acc_delete on nvidia targets. */ + +/* { dg-do run { target openacc_nvidia_accel_selected } } */ #include #include diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-28.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-28.c index 9304daa..7a96ab2 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-28.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-28.c @@ -1,4 +1,6 @@ -/* { dg-do run } */ +/* Exercise acc_delete with a NULL address on nvidia targets. */ + +/* { dg-do run { target openacc_nvidia_accel_selected } } */ #include #include diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-29.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-29.c index 92e3858..318a060 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-29.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-29.c @@ -1,4 +1,6 @@ -/* { dg-do run } */ +/* Exercise acc_delete with size zero on nvidia targets. */ + +/* { dg-do run { target openacc_nvidia_accel_selected } } */ #include #include diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-30.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-30.c index e81627d..9bc9ecc 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-30.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-30.c @@ -1,4 +1,6 @@ -/* { dg-do run } */ +/* Exercise an invalid partial acc_delete on nvidia targets. */ + +/* { dg-do run { target openacc_nvidia_accel_selected } } */ #include #include diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-34.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-34.c index 031c731..a24916d 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-34.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-34.c @@ -1,4 +1,6 @@ -/* { dg-do run } */ +/* Exercise an invalid acc_present_or_create on nvidia targets. */ + +/* { dg-do run { target openacc_nvidia_accel_selected } } */ #include #include diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-42.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-42.c index de5d1c1..30b90d4 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-42.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-42.c @@ -1,4 +1,6 @@ -/* { dg-do run } */ +/* Exercise acc_update_device on unmapped data on nvidia targets. */ + +/* { dg-do run { target openacc_nvidia_accel_selected } } */ #include #include diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-43.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-43.c index 0d593f0..5db2912 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-43.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-43.c @@ -1,4 +1,6 @@ -/* { dg-do run } */ +/* Exercise acc_update_device with a NULL data address on nvidia targets. */ + +/* { dg-do run { target openacc_nvidia_accel_selected } } */ #include #include diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-44.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-44.c index e98ecc4..8bbf016 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-44.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-44.c @@ -1,4 +1,6 @@ -/* { dg-do run } */ +/* Exercise acc_update_device with size zero data on nvidia targets. */ + +/* { dg-do run { target openacc_nvidia_accel_selected } } */ #include #include diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-47.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-47.c index f26fc33..c214042 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-47.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-47.c @@ -1,4 +1,6 @@ -/* { dg-do run } */ +/* Exercise acc_update_self with a NULL data mapping on nvidia targets. */ + +/* { dg-do run { target openacc_nvidia_accel_selected } } */ #include #include diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-48.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-48.c index 253ce59..afa137f 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-48.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-48.c @@ -1,4 +1,6 @@ -/* { dg-do run } */ +/* Exercise acc_update_self with a size zero data mapping on nvidia targets. */ + +/* { dg-do run { target openacc_nvidia_accel_selected } } */ #include #include diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-52.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-52.c index cfbb077..25c70c2 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-52.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-52.c @@ -1,4 +1,6 @@ -/* { dg-do run } */ +/* Exercise acc_map_data with a NULL data mapping on nvidia targets. */ + +/* { dg-do run { target openacc_nvidia_accel_selected } } */ #include #include diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-53.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-53.c index 5de376d..a8ee7df 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-53.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-53.c @@ -1,4 +1,6 @@ -/* { dg-do run } */ +/* Exercise acc_map_data with a NULL data mapping on nvidia targets. */ + +/* { dg-do run { target openacc_nvidia_accel_selected } } */ #include #include diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-54.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-54.c index 3e621c3..fc221f4 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-54.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-54.c @@ -1,4 +1,6 @@ -/* { dg-do run } */ +/* Exercise acc_map_data with data size of zero on nvidia targets. */ + +/* { dg-do run { target openacc_nvidia_accel_selected } } */ #include #include