From patchwork Tue Feb 9 15:00:31 2016 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Cesar Philippidis X-Patchwork-Id: 580883 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 53E8C140784 for ; Wed, 10 Feb 2016 02:01:09 +1100 (AEDT) Authentication-Results: ozlabs.org; dkim=pass (1024-bit key; unprotected) header.d=gcc.gnu.org header.i=@gcc.gnu.org header.b=HDh4gkBN; 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:from:message-id:date:mime-version :in-reply-to:content-type; q=dns; s=default; b=jkMQjtlpdZv+Dov1P uxe06f5l0Y/R9JBoSDLvrq7JAQxSq3Zlu32B9hxI5gfjff1UQeUM5aI67KD3pVUf 7ULVPDgoC4gDWi8ipvFR5UIyglsFgHOUbln6ImP/4zPYNnUJVYw7TOtsXDihce3q GlijrMYFdKqV1LEIZ9GnrXBWEA= 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:from:message-id:date:mime-version :in-reply-to:content-type; s=default; bh=xLDSgBdAGG1Ah6vg36UBLx9 BFcw=; b=HDh4gkBN2p52YSyXG29zx2bJpAhOeoAiVcU8jnEcqdS5srgyMGX9fij 2nxpZWo5eA9UrQVc6x9TyDbw4B/A2AtmkRJEbVTl14i9K//bU4YAZ6oyACc/tWY6 SA4Ir/JGBdsY5JSqhD67fbq0O6XPKXRQO2RIpZDJAvjiFAbC8QKc= Received: (qmail 100388 invoked by alias); 9 Feb 2016 15:00:53 -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 100230 invoked by uid 89); 9 Feb 2016 15:00:45 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=0.8 required=5.0 tests=BAYES_50, RCVD_IN_DNSWL_NONE, SPF_PASS autolearn=ham version=3.3.2 spammy=sk:update-, sk:update, 811, undeclared X-HELO: relay1.mentorg.com Received: from relay1.mentorg.com (HELO relay1.mentorg.com) (192.94.38.131) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with ESMTP; Tue, 09 Feb 2016 15:00:37 +0000 Received: from svr-orw-fem-06.mgc.mentorg.com ([147.34.97.120]) by relay1.mentorg.com with esmtp id 1aT9mK-00058w-IL from Cesar_Philippidis@mentor.com ; Tue, 09 Feb 2016 07:00:32 -0800 Received: from [127.0.0.1] (147.34.91.1) by SVR-ORW-FEM-06.mgc.mentorg.com (147.34.97.120) with Microsoft SMTP Server id 14.3.224.2; Tue, 9 Feb 2016 07:00:32 -0800 Subject: Re: [openacc] reference-typed data mappings To: "gcc-patches@gcc.gnu.org" , Jakub Jelinek References: <56AF9C98.4090200@codesourcery.com> From: Cesar Philippidis Message-ID: <56B9FF0F.6090904@mentor.com> Date: Tue, 9 Feb 2016 07:00:31 -0800 User-Agent: Mozilla/5.0 (X11; Linux x86_64; rv:38.0) Gecko/20100101 Thunderbird/38.5.1 MIME-Version: 1.0 In-Reply-To: <56AF9C98.4090200@codesourcery.com> On 02/01/2016 09:57 AM, Cesar Philippidis wrote: > This patch fixes a couple of bugs preventing c++ reference-typed > variables from working in openacc data clauses. These fixes include: > > * Teach the gimplifier to filter out pointer data mappings for > OACC_DATA, OACC_ENTER_DATA, OACC_EXIT_DATA and OACC_UPDATE regions. > Along with using a firsptrivate mapping for the array base pointers > in OACC_DATA, OACC_PARALLEL and OACC_KERNELS regions. > > * Make the data mapping errors emitted by the c and c++ front ends > more consistent with openacc by reporting data mapping errors, not > omp-specific map errors. > > * Add some light checking for duplicate reference mappings in c++. The > c++ FE still fails to detect duplicate component refs, but that's not > working in openacc at the moment, anyway. > > Jakub, the latter issue also affects openmp. I've added a simple openmp > test case, but it could probably be more extensive. Can you add more > test coverage or tell me what should be included? While working on a different reduction problem, I noticed that both the c and c++ front end's are treating reductions as generic data clauses. That means, parallel reductions of the form #pragma acc copy(foo) reduction(+:foo) would get treated as an error. This patch fixes that, in addition to the changes listed above. Is this patch ok for trunk? Cesar 2016-02-09 Cesar Philippidis gcc/c/ * c-parser.c (c_parser_oacc_all_clauses): Update call to c_finish_omp_clauses. (c_parser_omp_all_clauses): Likewise. (c_parser_oacc_cache): Likewise. (c_parser_oacc_loop): Likewise. (omp_split_clauses): Likewise. (c_parser_omp_declare_target): Likewise. (c_parser_cilk_for): Likewise. * c-tree.h (c_finish_omp_clauses): Update prototype. * c-typeck.c (c_finish_omp_clauses): Add is_oacc argument. Report OMP_CLAUSE_MAP errors as data clause errors for OpenACC. Allow OpenACC reductions variables to appear in data clauses. gcc/cp/ * cp-tree.h (finish_omp_clauses): Update prototype. * parser.c (cp_parser_oacc_all_clauses): Update call to finish_omp_clauses. (cp_parser_omp_all_clauses): Likewise. (cp_parser_omp_for_loop): Likewise. (cp_omp_split_clauses): Likewise. (cp_parser_oacc_cache): Likewise. (cp_parser_oacc_loop): Likewise. (cp_parser_omp_declare_target): Likewise. (cp_parser_cilk_for): Likewise. * pt.c (tsubst_attribute): Update call to tsubst_omp_clauses. (tsubst_omp_clauses): New is_oacc argument. Use it when calling finish_omp_clauses. (tsubst_omp_for_iterator): Update call to finish_omp_clauses. (tsubst_expr): Update calls to tsubst_omp_clauses. * semantics.c (finish_omp_clauses): Add is_oacc argument. Report OMP_CLAUSE_MAP errors as data clause errors for OpenACC. Check for duplicate reference mappings. Exclude "omp declare simd"-isms when processing OpenACC clauses. Allow OpenACC reductions variables to appear in data clauses. (finish_omp_for): Update call to finish_omp_clauses. gcc/ * gimplify.c (gimplify_scan_omp_clauses): Consider OACC_{DATA, PARALLEL, KERNELS} when setting target_firstprivatize_array_bases. Consider OACC_{DATA, ENTER_DATA, EXIT_DATA, UPDATE} when filtering out pointer mappings. Also filter out GOMP_MAP_POINTER. gcc/testsuite/ * c-c++-common/goacc/data-clause-duplicate-1.c: Adjust test. * c-c++-common/goacc/declare-2.c: Likewise. * c-c++-common/goacc/deviceptr-1.c: Likewise. * c-c++-common/goacc/kernels-alias-ipa-pta-3.c: Likewise. * c-c++-common/goacc/parallel-reduction.c: New test. * c-c++-common/goacc/pcopy.c: Adjust test. * 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/present-1.c: Likewise. * c-c++-common/goacc/private-reduction-1.c: New test. * c-c++-common/goacc/reduction-5.c: New test. * g++.dg/goacc/data-1.C: New test. * g++.dg/goacc/data-2.C: New test. * g++.dg/goacc/reduction-1.C: New test. * g++.dg/goacc/update.C: New test. * g++.dg/gomp/template-data.C: New test. libgomp/ * testsuite/libgomp.c++/non-scalar-data.C: New test. * testsuite/libgomp.oacc-c++/data-references.C: New test. * testsuite/libgomp.oacc-c++/data-templates.C: New test. * testsuite/libgomp.oacc-c++/non-scalar-data-templates.C: New test. * testsuite/libgomp.oacc-c++/update-reference.C: New test. * testsuite/libgomp.oacc-c++/update-template.C: New test. * testsuite/libgomp.oacc-c-c++-common/asyncwait-1.c: Adjust test. * testsuite/libgomp.oacc-c-c++-common/data-2.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/data-3.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/if-1.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/nested-1.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/present-2.c: Likewise. diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c index eede3a7..20ff7da 100644 --- a/gcc/c/c-parser.c +++ b/gcc/c/c-parser.c @@ -13115,7 +13115,7 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask, c_parser_skip_to_pragma_eol (parser); if (finish_p) - return c_finish_omp_clauses (clauses, false); + return c_finish_omp_clauses (clauses, false, true, false); return clauses; } @@ -13400,8 +13400,8 @@ c_parser_omp_all_clauses (c_parser *parser, omp_clause_mask mask, if (finish_p) { if ((mask & (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_UNIFORM)) != 0) - return c_finish_omp_clauses (clauses, true, true); - return c_finish_omp_clauses (clauses, true); + return c_finish_omp_clauses (clauses, true, false, true); + return c_finish_omp_clauses (clauses, true, false); } return clauses; @@ -13435,7 +13435,7 @@ c_parser_oacc_cache (location_t loc, c_parser *parser) tree stmt, clauses; clauses = c_parser_omp_var_list_parens (parser, OMP_CLAUSE__CACHE_, NULL); - clauses = c_finish_omp_clauses (clauses, false); + clauses = c_finish_omp_clauses (clauses, false, true); c_parser_skip_to_pragma_eol (parser); @@ -13769,9 +13769,9 @@ c_parser_oacc_loop (location_t loc, c_parser *parser, char *p_name, { clauses = c_oacc_split_loop_clauses (clauses, cclauses); if (*cclauses) - c_finish_omp_clauses (*cclauses, false); + c_finish_omp_clauses (*cclauses, false, true); if (clauses) - c_finish_omp_clauses (clauses, false); + c_finish_omp_clauses (clauses, false, true); } tree block = c_begin_compound_stmt (true); @@ -14942,7 +14942,7 @@ omp_split_clauses (location_t loc, enum tree_code code, c_omp_split_clauses (loc, code, mask, clauses, cclauses); for (i = 0; i < C_OMP_CLAUSE_SPLIT_COUNT; i++) if (cclauses[i]) - cclauses[i] = c_finish_omp_clauses (cclauses[i], true); + cclauses[i] = c_finish_omp_clauses (cclauses[i], true, false); } /* OpenMP 4.0: @@ -16455,7 +16455,7 @@ c_parser_omp_declare_target (c_parser *parser) { clauses = c_parser_omp_var_list_parens (parser, OMP_CLAUSE_TO_DECLARE, clauses); - clauses = c_finish_omp_clauses (clauses, true); + clauses = c_finish_omp_clauses (clauses, true, false); c_parser_skip_to_pragma_eol (parser); } else @@ -17503,7 +17503,7 @@ c_parser_cilk_for (c_parser *parser, tree grain) tree clauses = build_omp_clause (EXPR_LOCATION (grain), OMP_CLAUSE_SCHEDULE); OMP_CLAUSE_SCHEDULE_KIND (clauses) = OMP_CLAUSE_SCHEDULE_CILKFOR; OMP_CLAUSE_SCHEDULE_CHUNK_EXPR (clauses) = grain; - clauses = c_finish_omp_clauses (clauses, false); + clauses = c_finish_omp_clauses (clauses, false, false); tree block = c_begin_compound_stmt (true); tree sb = push_stmt_list (); @@ -17568,7 +17568,7 @@ c_parser_cilk_for (c_parser *parser, tree grain) OMP_CLAUSE_OPERAND (c, 0) = cilk_for_number_of_iterations (omp_for); OMP_CLAUSE_CHAIN (c) = clauses; - OMP_PARALLEL_CLAUSES (omp_par) = c_finish_omp_clauses (c, true); + OMP_PARALLEL_CLAUSES (omp_par) = c_finish_omp_clauses (c, true, false); add_stmt (omp_par); } diff --git a/gcc/c/c-tree.h b/gcc/c/c-tree.h index cf79ba7..152e0d1 100644 --- a/gcc/c/c-tree.h +++ b/gcc/c/c-tree.h @@ -660,7 +660,7 @@ extern tree c_begin_omp_task (void); extern tree c_finish_omp_task (location_t, tree, tree); extern void c_finish_omp_cancel (location_t, tree); extern void c_finish_omp_cancellation_point (location_t, tree); -extern tree c_finish_omp_clauses (tree, bool, bool = false); +extern tree c_finish_omp_clauses (tree, bool, bool, bool = false); extern tree c_build_va_arg (location_t, tree, location_t, tree); extern tree c_finish_transaction (location_t, tree, int); extern bool c_tree_equal (tree, tree); diff --git a/gcc/c/c-typeck.c b/gcc/c/c-typeck.c index 65925cb..149cdd5 100644 --- a/gcc/c/c-typeck.c +++ b/gcc/c/c-typeck.c @@ -12526,10 +12526,10 @@ c_find_omp_placeholder_r (tree *tp, int *, void *data) Remove any elements from the list that are invalid. */ tree -c_finish_omp_clauses (tree clauses, bool is_omp, bool declare_simd) +c_finish_omp_clauses (tree clauses, bool is_omp, bool is_oacc, bool declare_simd) { 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; @@ -12546,6 +12546,7 @@ c_finish_omp_clauses (tree clauses, bool is_omp, bool declare_simd) 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) { @@ -12866,6 +12867,16 @@ c_finish_omp_clauses (tree clauses, bool is_omp, bool declare_simd) omp_clause_code_name[OMP_CLAUSE_CODE (c)]); remove = true; } + else if (is_oacc && 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))) @@ -13157,8 +13168,10 @@ c_finish_omp_clauses (tree clauses, bool is_omp, bool declare_simd) { if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP) error ("%qD appears more than once in motion clauses", t); - else + else if (is_omp) error ("%qD appears more than once in map clauses", t); + else + error ("%qD appears more than once in data clauses", t); remove = true; } else if (bitmap_bit_p (&generic_head, DECL_UID (t)) diff --git a/gcc/cp/cp-tree.h b/gcc/cp/cp-tree.h index ead017e..2b70c96 100644 --- a/gcc/cp/cp-tree.h +++ b/gcc/cp/cp-tree.h @@ -6417,7 +6417,7 @@ extern tree omp_reduction_id (enum tree_code, tree, tree); extern tree cp_remove_omp_priv_cleanup_stmt (tree *, int *, void *); extern void cp_check_omp_declare_reduction (tree); extern void finish_omp_declare_simd_methods (tree); -extern tree finish_omp_clauses (tree, bool, bool = false); +extern tree finish_omp_clauses (tree, bool, bool, bool = false); extern tree push_omp_privatization_clauses (bool); extern void pop_omp_privatization_clauses (tree); extern void save_omp_privatization_clauses (vec &); diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c index d03b0c9..458abbe 100644 --- a/gcc/cp/parser.c +++ b/gcc/cp/parser.c @@ -32168,7 +32168,7 @@ cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask, cp_parser_skip_to_pragma_eol (parser, pragma_tok); if (finish_p) - return finish_omp_clauses (clauses, false); + return finish_omp_clauses (clauses, true, true); return clauses; } @@ -32487,9 +32487,9 @@ cp_parser_omp_all_clauses (cp_parser *parser, omp_clause_mask mask, if (finish_p) { if ((mask & (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_UNIFORM)) != 0) - return finish_omp_clauses (clauses, false, true); + return finish_omp_clauses (clauses, false, false, true); else - return finish_omp_clauses (clauses, true); + return finish_omp_clauses (clauses, false, true); } return clauses; } @@ -33564,7 +33564,7 @@ cp_parser_omp_for_loop (cp_parser *parser, enum tree_code code, tree clauses, else c = build_omp_clause (loc, OMP_CLAUSE_LASTPRIVATE); OMP_CLAUSE_DECL (c) = add_private_clause; - c = finish_omp_clauses (c, true); + c = finish_omp_clauses (c, false, true); if (c) { OMP_CLAUSE_CHAIN (c) = clauses; @@ -33713,7 +33713,7 @@ cp_omp_split_clauses (location_t loc, enum tree_code code, c_omp_split_clauses (loc, code, mask, clauses, cclauses); for (i = 0; i < C_OMP_CLAUSE_SPLIT_COUNT; i++) if (cclauses[i]) - cclauses[i] = finish_omp_clauses (cclauses[i], true); + cclauses[i] = finish_omp_clauses (cclauses[i], false, true); } /* OpenMP 4.0: @@ -34985,7 +34985,7 @@ cp_parser_oacc_cache (cp_parser *parser, cp_token *pragma_tok) tree stmt, clauses; clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE__CACHE_, NULL_TREE); - clauses = finish_omp_clauses (clauses, false); + clauses = finish_omp_clauses (clauses, true, true); cp_parser_require_pragma_eol (parser, cp_lexer_peek_token (parser->lexer)); @@ -35310,9 +35310,9 @@ cp_parser_oacc_loop (cp_parser *parser, cp_token *pragma_tok, char *p_name, { clauses = c_oacc_split_loop_clauses (clauses, cclauses); if (*cclauses) - finish_omp_clauses (*cclauses, false); + finish_omp_clauses (*cclauses, true, true); if (clauses) - finish_omp_clauses (clauses, false); + finish_omp_clauses (clauses, true, true); } tree block = begin_omp_structured_block (); @@ -35678,7 +35678,7 @@ cp_parser_omp_declare_target (cp_parser *parser, cp_token *pragma_tok) { clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE_TO_DECLARE, clauses); - clauses = finish_omp_clauses (clauses, true); + clauses = finish_omp_clauses (clauses, false, true); cp_parser_require_pragma_eol (parser, pragma_tok); } else @@ -37652,7 +37652,7 @@ cp_parser_cilk_for (cp_parser *parser, tree grain) tree clauses = build_omp_clause (EXPR_LOCATION (grain), OMP_CLAUSE_SCHEDULE); OMP_CLAUSE_SCHEDULE_KIND (clauses) = OMP_CLAUSE_SCHEDULE_CILKFOR; OMP_CLAUSE_SCHEDULE_CHUNK_EXPR (clauses) = grain; - clauses = finish_omp_clauses (clauses, false); + clauses = finish_omp_clauses (clauses, false, false); tree ret = cp_parser_omp_for_loop (parser, CILK_FOR, clauses, NULL); if (ret) diff --git a/gcc/cp/pt.c b/gcc/cp/pt.c index 76a6019..278c110 100644 --- a/gcc/cp/pt.c +++ b/gcc/cp/pt.c @@ -9537,7 +9537,8 @@ can_complete_type_without_circularity (tree type) return 1; } -static tree tsubst_omp_clauses (tree, bool, bool, tree, tsubst_flags_t, tree); +static tree tsubst_omp_clauses (tree, bool, bool, tree, tsubst_flags_t, tree, + bool); /* Instantiate a single dependent attribute T (a TREE_LIST), and return either T or a new TREE_LIST, possibly a chain in the case of a pack expansion. */ @@ -9557,9 +9558,9 @@ tsubst_attribute (tree t, tree *decl_p, tree args, { tree clauses = TREE_VALUE (val); clauses = tsubst_omp_clauses (clauses, true, false, args, - complain, in_decl); + complain, in_decl, false); c_omp_declare_simd_clauses_to_decls (*decl_p, clauses); - clauses = finish_omp_clauses (clauses, false, true); + clauses = finish_omp_clauses (clauses, false, false, true); tree parms = DECL_ARGUMENTS (*decl_p); clauses = c_omp_declare_simd_clauses_to_numbers (parms, clauses); @@ -14456,7 +14457,8 @@ tsubst_omp_clause_decl (tree decl, tree args, tsubst_flags_t complain, static tree tsubst_omp_clauses (tree clauses, bool declare_simd, bool allow_fields, - tree args, tsubst_flags_t complain, tree in_decl) + tree args, tsubst_flags_t complain, tree in_decl, + bool is_oacc) { tree new_clauses = NULL_TREE, nc, oc; tree linear_no_step = NULL_TREE; @@ -14669,7 +14671,7 @@ tsubst_omp_clauses (tree clauses, bool declare_simd, bool allow_fields, new_clauses = nreverse (new_clauses); if (!declare_simd) { - new_clauses = finish_omp_clauses (new_clauses, allow_fields); + new_clauses = finish_omp_clauses (new_clauses, is_oacc, allow_fields); if (linear_no_step) for (nc = new_clauses; nc; nc = OMP_CLAUSE_CHAIN (nc)) if (nc == linear_no_step) @@ -14890,7 +14892,7 @@ tsubst_omp_for_iterator (tree t, int i, tree declv, tree orig_declv, { tree c = build_omp_clause (input_location, OMP_CLAUSE_PRIVATE); OMP_CLAUSE_DECL (c) = decl; - c = finish_omp_clauses (c, true); + c = finish_omp_clauses (c, false, true); if (c) { OMP_CLAUSE_CHAIN (c) = *clauses; @@ -15373,8 +15375,8 @@ tsubst_expr (tree t, tree args, tsubst_flags_t complain, tree in_decl, case OACC_KERNELS: case OACC_PARALLEL: - tmp = tsubst_omp_clauses (OMP_CLAUSES (t), false, false, args, complain, - in_decl); + tmp = tsubst_omp_clauses (OMP_CLAUSES (t), false, true, args, complain, + in_decl, true); stmt = begin_omp_parallel (); RECUR (OMP_BODY (t)); finish_omp_construct (TREE_CODE (t), stmt, tmp); @@ -15383,7 +15385,7 @@ tsubst_expr (tree t, tree args, tsubst_flags_t complain, tree in_decl, case OMP_PARALLEL: r = push_omp_privatization_clauses (OMP_PARALLEL_COMBINED (t)); tmp = tsubst_omp_clauses (OMP_PARALLEL_CLAUSES (t), false, true, - args, complain, in_decl); + args, complain, in_decl, false); if (OMP_PARALLEL_COMBINED (t)) omp_parallel_combined_clauses = &tmp; stmt = begin_omp_parallel (); @@ -15397,7 +15399,7 @@ tsubst_expr (tree t, tree args, tsubst_flags_t complain, tree in_decl, case OMP_TASK: r = push_omp_privatization_clauses (false); tmp = tsubst_omp_clauses (OMP_TASK_CLAUSES (t), false, true, - args, complain, in_decl); + args, complain, in_decl, false); stmt = begin_omp_task (); RECUR (OMP_TASK_BODY (t)); finish_omp_task (tmp, stmt); @@ -15419,9 +15421,9 @@ tsubst_expr (tree t, tree args, tsubst_flags_t complain, tree in_decl, int i; r = push_omp_privatization_clauses (OMP_FOR_INIT (t) == NULL_TREE); - clauses = tsubst_omp_clauses (OMP_FOR_CLAUSES (t), false, - TREE_CODE (t) != OACC_LOOP, - args, complain, in_decl); + clauses = tsubst_omp_clauses (OMP_FOR_CLAUSES (t), false, true, + args, complain, in_decl, + TREE_CODE (t) == OACC_LOOP); if (OMP_FOR_INIT (t) != NULL_TREE) { declv = make_tree_vec (TREE_VEC_LENGTH (OMP_FOR_INIT (t))); @@ -15478,7 +15480,7 @@ tsubst_expr (tree t, tree args, tsubst_flags_t complain, tree in_decl, r = push_omp_privatization_clauses (TREE_CODE (t) == OMP_TEAMS && OMP_TEAMS_COMBINED (t)); tmp = tsubst_omp_clauses (OMP_CLAUSES (t), false, true, - args, complain, in_decl); + args, complain, in_decl, false); stmt = push_stmt_list (); RECUR (OMP_BODY (t)); stmt = pop_stmt_list (stmt); @@ -15493,9 +15495,8 @@ tsubst_expr (tree t, tree args, tsubst_flags_t complain, tree in_decl, case OACC_DATA: case OMP_TARGET_DATA: case OMP_TARGET: - tmp = tsubst_omp_clauses (OMP_CLAUSES (t), false, - TREE_CODE (t) != OACC_DATA, - args, complain, in_decl); + tmp = tsubst_omp_clauses (OMP_CLAUSES (t), false, true, args, complain, + in_decl, TREE_CODE (t) == OACC_DATA); keep_next_level (true); stmt = begin_omp_structured_block (); @@ -15540,8 +15541,8 @@ tsubst_expr (tree t, tree args, tsubst_flags_t complain, tree in_decl, case OACC_DECLARE: t = copy_node (t); - tmp = tsubst_omp_clauses (OACC_DECLARE_CLAUSES (t), false, false, - args, complain, in_decl); + tmp = tsubst_omp_clauses (OACC_DECLARE_CLAUSES (t), false, true, + args, complain, in_decl, true); OACC_DECLARE_CLAUSES (t) = tmp; add_stmt (t); break; @@ -15550,7 +15551,7 @@ tsubst_expr (tree t, tree args, tsubst_flags_t complain, tree in_decl, case OMP_TARGET_ENTER_DATA: case OMP_TARGET_EXIT_DATA: tmp = tsubst_omp_clauses (OMP_STANDALONE_CLAUSES (t), false, true, - args, complain, in_decl); + args, complain, in_decl, false); t = copy_node (t); OMP_STANDALONE_CLAUSES (t) = tmp; add_stmt (t); @@ -15559,8 +15560,8 @@ tsubst_expr (tree t, tree args, tsubst_flags_t complain, tree in_decl, case OACC_ENTER_DATA: case OACC_EXIT_DATA: case OACC_UPDATE: - tmp = tsubst_omp_clauses (OMP_STANDALONE_CLAUSES (t), false, false, - args, complain, in_decl); + tmp = tsubst_omp_clauses (OMP_STANDALONE_CLAUSES (t), false, true, + args, complain, in_decl, true); t = copy_node (t); OMP_STANDALONE_CLAUSES (t) = tmp; add_stmt (t); @@ -15568,7 +15569,7 @@ tsubst_expr (tree t, tree args, tsubst_flags_t complain, tree in_decl, case OMP_ORDERED: tmp = tsubst_omp_clauses (OMP_ORDERED_CLAUSES (t), false, true, - args, complain, in_decl); + args, complain, in_decl, false); stmt = push_stmt_list (); RECUR (OMP_BODY (t)); stmt = pop_stmt_list (stmt); diff --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c index c9f9db4..ae6ba0f 100644 --- a/gcc/cp/semantics.c +++ b/gcc/cp/semantics.c @@ -5736,10 +5736,11 @@ cp_finish_omp_clause_depend_sink (tree sink_clause) Remove any elements from the list that are invalid. */ tree -finish_omp_clauses (tree clauses, bool allow_fields, bool declare_simd) +finish_omp_clauses (tree clauses, bool is_oacc, bool allow_fields, + bool declare_simd) { 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; @@ -5753,6 +5754,7 @@ finish_omp_clauses (tree clauses, bool allow_fields, bool declare_simd) 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) { @@ -5966,6 +5968,16 @@ finish_omp_clauses (tree clauses, bool allow_fields, bool declare_simd) omp_clause_code_name[OMP_CLAUSE_CODE (c)]); remove = true; } + else if (is_oacc && 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))) @@ -5976,7 +5988,10 @@ finish_omp_clauses (tree clauses, bool allow_fields, bool declare_simd) 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 (is_oacc) + error ("%qD appears more than once in data clauses", t); + else + error ("%qD appears both in data and map clauses", t); remove = true; } else @@ -6002,7 +6017,8 @@ finish_omp_clauses (tree clauses, bool allow_fields, bool declare_simd) omp_note_field_privatization (t, OMP_CLAUSE_DECL (c)); else t = OMP_CLAUSE_DECL (c); - if (t == current_class_ptr) + if (!is_oacc + && t == current_class_ptr) { error ("% allowed in OpenMP only in %" " clauses"); @@ -6028,7 +6044,10 @@ finish_omp_clauses (tree clauses, bool allow_fields, bool declare_simd) } else if (bitmap_bit_p (&map_head, DECL_UID (t))) { - error ("%qD appears both in data and map clauses", t); + if (is_oacc) + error ("%qD appears more than once in data clauses", t); + else + error ("%qD appears both in data and map clauses", t); remove = true; } else @@ -6538,6 +6557,9 @@ finish_omp_clauses (tree clauses, bool allow_fields, bool declare_simd) if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP) error ("%qD appears more than once in motion" " clauses", t); + else if (is_oacc) + error ("%qD appears more than once in data" + " clauses", t); else error ("%qD appears more than once in map" " clauses", t); @@ -6549,6 +6571,27 @@ finish_omp_clauses (tree clauses, bool allow_fields, bool declare_simd) bitmap_set_bit (&map_field_head, DECL_UID (t)); } } + else if (TREE_CODE (t) == TREE_LIST) + { + while (TREE_CODE (t = TREE_CHAIN (t)) == TREE_LIST) + ; + + if (DECL_P (t)) + { + if (bitmap_bit_p (&map_head, DECL_UID (t))) + { + if (is_oacc) + error ("%qD appears more than once in data " + "clauses", t); + else + error ("%qD appears more than once in map " + "clauses", t); + remove = true; + } + else + bitmap_set_bit (&map_head, DECL_UID (t)); + } + } } break; } @@ -6625,7 +6668,8 @@ finish_omp_clauses (tree clauses, bool allow_fields, bool declare_simd) omp_clause_code_name[OMP_CLAUSE_CODE (c)]); remove = true; } - else if (t == current_class_ptr) + else if (!is_oacc + && t == current_class_ptr) { error ("% allowed in OpenMP only in %" " clauses"); @@ -6666,7 +6710,10 @@ finish_omp_clauses (tree clauses, bool allow_fields, bool declare_simd) } else if (bitmap_bit_p (&map_head, DECL_UID (t))) { - error ("%qD appears both in data and map clauses", t); + if (is_oacc) + error ("%qD appears more than once in data clauses", t); + else + error ("%qD appears both in data and map clauses", t); remove = true; } else @@ -6676,6 +6723,8 @@ finish_omp_clauses (tree clauses, bool allow_fields, bool declare_simd) { if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP) error ("%qD appears more than once in motion clauses", t); + if (is_oacc) + error ("%qD appears more than once in data clauses", t); else error ("%qD appears more than once in map clauses", t); remove = true; @@ -6683,7 +6732,10 @@ finish_omp_clauses (tree clauses, bool allow_fields, bool declare_simd) 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 (is_oacc) + error ("%qD appears more than once in data clauses", t); + else + error ("%qD appears both in data and map clauses", t); remove = true; } else @@ -8260,7 +8312,7 @@ finish_omp_for (location_t locus, enum tree_code code, tree declv, OMP_CLAUSE_OPERAND (c, 0) = cilk_for_number_of_iterations (omp_for); OMP_CLAUSE_CHAIN (c) = clauses; - OMP_PARALLEL_CLAUSES (omp_par) = finish_omp_clauses (c, false); + OMP_PARALLEL_CLAUSES (omp_par) = finish_omp_clauses (c, false, false); add_stmt (omp_par); return omp_par; } diff --git a/gcc/gimplify.c b/gcc/gimplify.c index b0ee27e..6584bb7 100644 --- a/gcc/gimplify.c +++ b/gcc/gimplify.c @@ -6495,7 +6495,10 @@ 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_DATA: case OACC_HOST_DATA: + case OACC_PARALLEL: + case OACC_KERNELS: ctx->target_firstprivatize_array_bases = true; default: break; @@ -6761,10 +6764,16 @@ 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_DATA: case OACC_HOST_DATA: + case OACC_ENTER_DATA: + case OACC_EXIT_DATA: + case OACC_UPDATE: if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER || (OMP_CLAUSE_MAP_KIND (c) - == GOMP_MAP_FIRSTPRIVATE_REFERENCE)) + == GOMP_MAP_FIRSTPRIVATE_REFERENCE) + || (!lang_GNU_Fortran () && + OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER)) /* For target {,enter ,exit }data only the array slice is mapped, but not the pointer to it. */ remove = true; 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/declare-2.c b/gcc/testsuite/c-c++-common/goacc/declare-2.c index d24cb22..1b1e60a 100644 --- a/gcc/testsuite/c-c++-common/goacc/declare-2.c +++ b/gcc/testsuite/c-c++-common/goacc/declare-2.c @@ -77,3 +77,5 @@ f (void) #pragma acc declare present (v9) /* { dg-error "invalid use of" } */ } + +/* { dg-error "at file scope" "" { target c++ } 10 } */ diff --git a/gcc/testsuite/c-c++-common/goacc/deviceptr-1.c b/gcc/testsuite/c-c++-common/goacc/deviceptr-1.c index 546fa82..bc10b82 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-ipa-pta-3.c b/gcc/testsuite/c-c++-common/goacc/kernels-alias-ipa-pta-3.c index 1eb56eb..97d8f52 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,6 @@ 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 } } } */ +/* { dg-final { scan-tree-dump-times "(?n)= \\*a" 1 "optimized" { target c } } } */ diff --git a/gcc/testsuite/c-c++-common/goacc/parallel-reduction.c b/gcc/testsuite/c-c++-common/goacc/parallel-reduction.c new file mode 100644 index 0000000..d7cc947 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/parallel-reduction.c @@ -0,0 +1,17 @@ +int +main () +{ + 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; +} diff --git a/gcc/testsuite/c-c++-common/goacc/pcopy.c b/gcc/testsuite/c-c++-common/goacc/pcopy.c index 02c4383..7e39e42 100644 --- a/gcc/testsuite/c-c++-common/goacc/pcopy.c +++ b/gcc/testsuite/c-c++-common/goacc/pcopy.c @@ -7,4 +7,6 @@ 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\\(alloc:cp \\\[pointer assign, bias: 3]\\)" 1 "original" { target c } } } */ + +/* { dg-final { scan-tree-dump-times "#pragma acc parallel map\\(tofrom:\\*\\(cp \\+ 3\\) \\\[len: 5]\\) map\\(firstprivate:cp \\\[pointer assign, bias: 3]\\)" 1 "original" { target c++ } } } */ diff --git a/gcc/testsuite/c-c++-common/goacc/pcopyin.c b/gcc/testsuite/c-c++-common/goacc/pcopyin.c index 10911fc..f2c0354 100644 --- a/gcc/testsuite/c-c++-common/goacc/pcopyin.c +++ b/gcc/testsuite/c-c++-common/goacc/pcopyin.c @@ -7,4 +7,6 @@ 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\\(alloc:cp \\\[pointer assign, bias: 4]\\)" 1 "original" { target c } } } */ + +/* { dg-final { scan-tree-dump-times "#pragma acc parallel map\\(to:\\*\\(cp \\+ 4\\) \\\[len: 6]\\) map\\(firstprivate:cp \\\[pointer assign, bias: 4]\\)" 1 "original" { target c++ } } } */ diff --git a/gcc/testsuite/c-c++-common/goacc/pcopyout.c b/gcc/testsuite/c-c++-common/goacc/pcopyout.c index 703ac2f..555a0b6 100644 --- a/gcc/testsuite/c-c++-common/goacc/pcopyout.c +++ b/gcc/testsuite/c-c++-common/goacc/pcopyout.c @@ -7,4 +7,6 @@ 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\\(alloc:cp \\\[pointer assign, bias: 5]\\)" 1 "original" { target c } } } */ + +/* { dg-final { scan-tree-dump-times "#pragma acc parallel map\\(from:\\*\\(cp \\+ 5\\) \\\[len: 7]\\) map\\(firstprivate:cp \\\[pointer assign, bias: 5]\\)" 1 "original" { target c++ } } } */ diff --git a/gcc/testsuite/c-c++-common/goacc/pcreate.c b/gcc/testsuite/c-c++-common/goacc/pcreate.c index 00bf155..fa1b985 100644 --- a/gcc/testsuite/c-c++-common/goacc/pcreate.c +++ b/gcc/testsuite/c-c++-common/goacc/pcreate.c @@ -7,4 +7,6 @@ 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\\(alloc:cp \\\[pointer assign, bias: 6]\\)" 1 "original" { target c } } } */ + +/* { dg-final { scan-tree-dump-times "#pragma acc parallel map\\(alloc:\\*\\(cp \\+ 6\\) \\\[len: 8]\\) map\\(firstprivate:cp \\\[pointer assign, bias: 6]\\)" 1 "original" { target c++ } } } */ diff --git a/gcc/testsuite/c-c++-common/goacc/present-1.c b/gcc/testsuite/c-c++-common/goacc/present-1.c index 7537948..52ff4c6 100644 --- a/gcc/testsuite/c-c++-common/goacc/present-1.c +++ b/gcc/testsuite/c-c++-common/goacc/present-1.c @@ -7,4 +7,6 @@ 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\\(alloc:cp \\\[pointer assign, bias: 7]\\)" 1 "original" { target c } } } */ + +/* { 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" { target c++ } } } */ diff --git a/gcc/testsuite/c-c++-common/goacc/private-reduction-1.c b/gcc/testsuite/c-c++-common/goacc/private-reduction-1.c new file mode 100644 index 0000000..d4e3995 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/private-reduction-1.c @@ -0,0 +1,12 @@ +int +reduction () +{ + int i, r; + + #pragma acc parallel + #pragma acc loop private (r) reduction (+:r) + for (i = 0; i < 100; i++) + r += 10; + + return r; +} diff --git a/gcc/testsuite/c-c++-common/goacc/reduction-5.c b/gcc/testsuite/c-c++-common/goacc/reduction-5.c new file mode 100644 index 0000000..42cf9c2 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/reduction-5.c @@ -0,0 +1,21 @@ +int +main () +{ + int sum = 0; + int dummy = 0; + +#pragma acc data copy (dummy) + { +#pragma acc parallel num_gangs (10) copy (sum) reduction (+:sum) reduction (+:sum) /* { dg-error "appears more than once in reduction clauses" } */ + { + int i; + int v = 5; +#pragma acc loop reduction (+:sum, sum) /* { dg-error "appears more than once in reduction clauses" } */ + for (i = 0; i < 100; i++) + sum += 10 + v; + } + } + + return sum; +} + 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/gcc/testsuite/g++.dg/goacc/data-2.C b/gcc/testsuite/g++.dg/goacc/data-2.C new file mode 100644 index 0000000..efa002d --- /dev/null +++ b/gcc/testsuite/g++.dg/goacc/data-2.C @@ -0,0 +1,30 @@ +void +fun (float (&fp)[100]) +{ + float *dptr = &fp[50]; + +#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 data clauses" } */ + ; +#pragma acc data create(fp[:10]) deviceptr(dptr) + ; +#pragma acc data create(fp) present(fp) /* { dg-error "'fp' appears more than once in data clauses" } */ + ; +} + +template +void +fun (T (&fp)[100]) +{ + T *dptr = &fp[50]; + +#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 data clauses" } */ + ; +#pragma acc data create(fp[:10]) deviceptr(dptr) + ; +#pragma acc data create(fp) present(fp) /* { dg-error "'fp' appears more than once in data clauses" } */ + ; +} diff --git a/gcc/testsuite/g++.dg/goacc/reduction-1.C b/gcc/testsuite/g++.dg/goacc/reduction-1.C new file mode 100644 index 0000000..de97125 --- /dev/null +++ b/gcc/testsuite/g++.dg/goacc/reduction-1.C @@ -0,0 +1,72 @@ +/* { dg-require-effective-target alloca } */ +/* Integer reductions. */ + +#define vl 32 + +int +main(void) +{ + const int n = 1000; + int i; + int result, array[n]; + int lresult; + + /* '+' reductions. */ +#pragma acc parallel vector_length (vl) +#pragma acc loop reduction (+:result) + for (i = 0; i < n; i++) + result += array[i]; + + /* '*' reductions. */ +#pragma acc parallel vector_length (vl) +#pragma acc loop reduction (*:result) + for (i = 0; i < n; i++) + result *= array[i]; + +// result = 0; +// vresult = 0; +// +// /* 'max' reductions. */ +// #pragma acc parallel vector_length (vl) +// #pragma acc loop reduction (+:result) +// for (i = 0; i < n; i++) +// result = result > array[i] ? result : array[i]; +// +// /* 'min' reductions. */ +// #pragma acc parallel vector_length (vl) +// #pragma acc loop reduction (+:result) +// for (i = 0; i < n; i++) +// result = result < array[i] ? result : array[i]; + + /* '&' reductions. */ +#pragma acc parallel vector_length (vl) +#pragma acc loop reduction (&:result) + for (i = 0; i < n; i++) + result &= array[i]; + + /* '|' reductions. */ +#pragma acc parallel vector_length (vl) +#pragma acc loop reduction (|:result) + for (i = 0; i < n; i++) + result |= array[i]; + + /* '^' reductions. */ +#pragma acc parallel vector_length (vl) +#pragma acc loop reduction (^:result) + for (i = 0; i < n; i++) + result ^= array[i]; + + /* '&&' reductions. */ +#pragma acc parallel vector_length (vl) +#pragma acc loop reduction (&&:lresult) + for (i = 0; i < n; i++) + lresult = lresult && (result > array[i]); + + /* '||' reductions. */ +#pragma acc parallel vector_length (vl) +#pragma acc loop reduction (||:lresult) + for (i = 0; i < n; i++) + lresult = lresult || (result > array[i]); + + return 0; +} diff --git a/gcc/testsuite/g++.dg/goacc/update.C b/gcc/testsuite/g++.dg/goacc/update.C new file mode 100644 index 0000000..18091c9 --- /dev/null +++ b/gcc/testsuite/g++.dg/goacc/update.C @@ -0,0 +1,22 @@ +void +f (int &i, int (&a)[10]) +{ +#pragma acc update device(i) +#pragma acc update host(i) +#pragma acc update self(i) +#pragma acc update device(a[1:3]) +#pragma acc update host(a[1:3]) +#pragma acc update self(a[1:3]) +} + +template +void +f (T &i, T (&a)[10]) +{ +#pragma acc update device(i) +#pragma acc update host(i) +#pragma acc update self(i) +#pragma acc update device(a[1:3]) +#pragma acc update host(a[1:3]) +#pragma acc update self(a[1:3]) +} diff --git a/gcc/testsuite/g++.dg/gomp/template-data.C b/gcc/testsuite/g++.dg/gomp/template-data.C new file mode 100644 index 0000000..0be14d4 --- /dev/null +++ b/gcc/testsuite/g++.dg/gomp/template-data.C @@ -0,0 +1,18 @@ +void +fun (float (&fp)[100]) +{ + float *dptr = &fp[50]; + +#pragma omp target data map(tofrom:fp[0:2], fp[0:2]) /* { dg-error "'fp' appears more than once in data clauses" } */ + ; +} + +template +void +fun (T (&fp)[100]) +{ + T *dptr = &fp[50]; + +#pragma omp target data map(tofrom:fp[0:2], fp[0:2]) /* { dg-error "'fp' appears more than once in map clauses" } */ + ; +} diff --git a/libgomp/testsuite/libgomp.c++/non-scalar-data.C b/libgomp/testsuite/libgomp.c++/non-scalar-data.C new file mode 100644 index 0000000..180e86f --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/non-scalar-data.C @@ -0,0 +1,109 @@ +// Ensure that a non-scalar dummy arguments which are implicitly used inside +// offloaded regions are properly mapped using present_or_copy. + +// { dg-do run } + +#include + +const int n = 100; + +struct data { + int v; +}; + +void +kernels_present (data &d, int &x) +{ +#pragma acc kernels present (d, x) default (none) + { + d.v = x; + } +} + +void +parallel_present (data &d, int &x) +{ +#pragma acc parallel present (d, x) default (none) + { + d.v = x; + } +} + +void +kernels_implicit (data &d, int &x) +{ +#pragma acc kernels + { + d.v = x; + } +} + +void +parallel_implicit (data &d, int &x) +{ +#pragma acc parallel + { + d.v = x; + } +} + +void +reference_data (data &d, int &x) +{ +#pragma acc data copy(d, x) + { + kernels_present (d, x); + +#pragma acc update host(d) + assert (d.v == x); + + x = 200; +#pragma acc update device(x) + + parallel_present (d, x); + } + + assert (d.v = x); + + x = 300; + kernels_implicit (d, x); + assert (d.v = x); + + x = 400; + parallel_implicit (d, x); + assert (d.v = x); +} + +int +main () +{ + data d; + int x = 100; + +#pragma acc data copy(d, x) + { + kernels_present (d, x); + +#pragma acc update host(d) + assert (d.v == x); + + x = 200; +#pragma acc update device(x) + + parallel_present (d, x); + } + + assert (d.v = x); + + x = 300; + kernels_implicit (d, x); + assert (d.v = x); + + x = 400; + parallel_implicit (d, x); + assert (d.v = x); + + reference_data (d, x); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-c++/data-references.C b/libgomp/testsuite/libgomp.oacc-c++/data-references.C new file mode 100644 index 0000000..1f2abf4 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c++/data-references.C @@ -0,0 +1,178 @@ +/* Test data mappings on variables with reference types. */ + +/* { dg-do run } */ + +#include + +void +test (int &N, float *(&a), float *(&b), float *(&c), float *(&d), float *(&e)) +{ + int i; + + for (i = 0; i < N; i++) + { + a[i] = 3.0; + b[i] = 0.0; + } + +#pragma acc enter data copyin (a[0:N]) copyin (b[0:N]) copyin (N) async +#pragma acc parallel async wait present (N, a[0:N], b[0:N]) +#pragma acc loop + for (i = 0; i < N; i++) + b[i] = a[i]; + +#pragma acc exit data copyout (a[0:N]) copyout (b[0:N]) wait async +#pragma acc wait + + for (i = 0; i < N; i++) + { + if (a[i] != 3.0) + abort (); + + if (b[i] != 3.0) + abort (); + } + + for (i = 0; i < N; i++) + { + a[i] = 2.0; + b[i] = 0.0; + } + +#pragma acc enter data copyin (a[0:N]) copyin (b[0:N]) copyin (N) async (1) +#pragma acc parallel async (1) present (N, a[0:N], b[0:N]) +#pragma acc loop + for (i = 0; i < N; i++) + b[i] = a[i]; + +#pragma acc update host (a[0:N], b[0:N]) wait (1) async (1) +#pragma acc wait (1) + + for (i = 0; i < N; i++) + { + if (a[i] != 2.0) + abort (); + + if (b[i] != 2.0) + abort (); + } + + for (i = 0; i < N; i++) + { + a[i] = 3.0; + b[i] = 0.0; + c[i] = 0.0; + d[i] = 0.0; + } + +#pragma acc update device (a[0:N], b[0:N]) +#pragma acc enter data copyin (c[0:N], d[0:N]) async (1) + +#pragma acc parallel present (N, 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 present (N, 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 present (N, 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]; + +#pragma acc exit data copyout (a[0:N], b[0:N], c[0:N], d[0:N]) wait (1, 2, 3) async (1) +#pragma acc wait (1) + + 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; + } + +#pragma acc enter data copyin (a[0:N], b[0:N], c[0:N], d[0:N], e[0:N]) async (1) + +#pragma acc parallel loop present (N, 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 loop present (N, 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 loop present (N, 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 loop present (N, 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], b[0:N], c[0:N], d[0:N], e[0:N]) wait (1, 2, 3, 4) async (1) delete (N) +#pragma acc wait (1) + + 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 (); + } +} + +int +main (int argc, char **argv) +{ + int N = 128; //1024 * 1024; + float *a, *b, *c, *d, *e; + 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); + + test (N, a, b, c, d, e); + + free (a); + free (b); + free (c); + free (d); + free (e); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-c++/data-templates.C b/libgomp/testsuite/libgomp.oacc-c++/data-templates.C new file mode 100644 index 0000000..34bd556 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c++/data-templates.C @@ -0,0 +1,179 @@ +/* Test data mappings on variables with template types. */ + +/* { dg-do run } */ + +#include + +template +void +test (T1 &N, T2 *(&a), T2 *(&b), T2 *(&c), T2 *(&d), T2 *(&e)) +{ + int i; + + for (i = 0; i < N; i++) + { + a[i] = 3.0; + b[i] = 0.0; + } + +#pragma acc enter data copyin (a[0:N]) copyin (b[0:N]) copyin (N) async +#pragma acc parallel async wait present (N, a[0:N], b[0:N]) +#pragma acc loop + for (i = 0; i < N; i++) + b[i] = a[i]; + +#pragma acc exit data copyout (a[0:N]) copyout (b[0:N]) wait async +#pragma acc wait + + for (i = 0; i < N; i++) + { + if (a[i] != 3.0) + abort (); + + if (b[i] != 3.0) + abort (); + } + + for (i = 0; i < N; i++) + { + a[i] = 2.0; + b[i] = 0.0; + } + +#pragma acc enter data copyin (a[0:N]) copyin (b[0:N]) copyin (N) async (1) +#pragma acc parallel async (1) present (N, a[0:N], b[0:N]) +#pragma acc loop + for (i = 0; i < N; i++) + b[i] = a[i]; + +#pragma acc update host (a[0:N], b[0:N]) wait (1) async (1) +#pragma acc wait (1) + + for (i = 0; i < N; i++) + { + if (a[i] != 2.0) + abort (); + + if (b[i] != 2.0) + abort (); + } + + for (i = 0; i < N; i++) + { + a[i] = 3.0; + b[i] = 0.0; + c[i] = 0.0; + d[i] = 0.0; + } + +#pragma acc update device (a[0:N], b[0:N]) +#pragma acc enter data copyin (c[0:N], d[0:N]) async (1) + +#pragma acc parallel present (N, 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 present (N, 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 present (N, 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]; + +#pragma acc exit data copyout (a[0:N], b[0:N], c[0:N], d[0:N]) wait (1, 2, 3) async (1) +#pragma acc wait (1) + + 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; + } + +#pragma acc enter data copyin (a[0:N], b[0:N], c[0:N], d[0:N], e[0:N]) async (1) + +#pragma acc parallel loop present (N, 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 loop present (N, 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 loop present (N, 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 loop present (N, 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], b[0:N], c[0:N], d[0:N], e[0:N]) wait (1, 2, 3, 4) async (1) delete (N) +#pragma acc wait (1) + + 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 (); + } +} + +int +main (int argc, char **argv) +{ + int N = 128; //1024 * 1024; + float *a, *b, *c, *d, *e; + 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); + + test (N, a, b, c, d, e); + + free (a); + free (b); + free (c); + free (d); + free (e); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-c++/non-scalar-data-templates.C b/libgomp/testsuite/libgomp.oacc-c++/non-scalar-data-templates.C new file mode 100644 index 0000000..03b2b01 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c++/non-scalar-data-templates.C @@ -0,0 +1,114 @@ +// Ensure that a non-scalar dummy arguments which are implicitly used inside +// offloaded regions are properly mapped using present_or_copy. + +// { dg-do run } + +#include + +const int n = 100; + +struct data { + int v; +}; + +template +void +kernels_present (t1 &d, t2 &x) +{ +#pragma acc kernels present (d, x) default (none) + { + d.v = x; + } +} + +template +void +parallel_present (t1 &d, t2 &x) +{ +#pragma acc parallel present (d, x) default (none) + { + d.v = x; + } +} + +template +void +kernels_implicit (t1 &d, t2 &x) +{ +#pragma acc kernels + { + d.v = x; + } +} + +template +void +parallel_implicit (t1 &d, t2 &x) +{ +#pragma acc parallel + { + d.v = x; + } +} + +template +void +reference_data (t1 &d, t2 &x) +{ +#pragma acc data copy(d, x) + { + kernels_present (d, x); + +#pragma acc update host(d) + assert (d.v == x); + + x = 200; +#pragma acc update device(x) + + parallel_present (d, x); + } + + assert (d.v = x); + + x = 300; + kernels_implicit (d, x); + assert (d.v = x); + + x = 400; + parallel_implicit (d, x); + assert (d.v = x); +} + +int +main () +{ + data d; + int x = 100; + +#pragma acc data copy(d, x) + { + kernels_present (d, x); + +#pragma acc update host(d) + assert (d.v == x); + + x = 200; +#pragma acc update device(x) + + parallel_present (d, x); + } + + assert (d.v = x); + + x = 300; + kernels_implicit (d, x); + assert (d.v = x); + + x = 400; + parallel_implicit (d, x); + assert (d.v = x); + + reference_data (d, x); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-c++/update-reference.C b/libgomp/testsuite/libgomp.oacc-c++/update-reference.C new file mode 100644 index 0000000..bd85d13 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c++/update-reference.C @@ -0,0 +1,300 @@ +/* Copy of update-1.c with self exchanged with host for #pragma acc update. + This exercises references types. */ + +/* { dg-do run } */ +/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */ + +#include +#include +#include +#include +#include + +int +test (int &N, float *(&a), float *(&b), float *(&c), float *(&d_a), + float *(&d_b), float *(&d_c), int &i) +{ + + + for (i = 0; i < N; i++) + { + a[i] = 3.0; + b[i] = 0.0; + } + + acc_map_data (a, d_a, N * sizeof (float)); + acc_map_data (b, d_b, N * sizeof (float)); + acc_map_data (c, d_c, N * sizeof (float)); + +#pragma acc update device (a[0:N], b[0:N]) + +#pragma acc parallel present (a[0:N], b[0:N]) + { + int ii; + + for (ii = 0; ii < N; ii++) + b[ii] = a[ii]; + } + +#pragma acc update self (a[0:N], b[0:N]) + + for (i = 0; i < N; i++) + { + if (a[i] != 3.0) + abort (); + + if (b[i] != 3.0) + abort (); + } + + if (!acc_is_present (&a[0], (N * sizeof (float)))) + abort (); + + if (!acc_is_present (&b[0], (N * sizeof (float)))) + abort (); + + for (i = 0; i < N; i++) + { + a[i] = 5.0; + b[i] = 1.0; + } + +#pragma acc update device (a[0:N], b[0:N]) + +#pragma acc parallel present (a[0:N], b[0:N]) + { + int ii; + + for (ii = 0; ii < N; ii++) + b[ii] = a[ii]; + } + +#pragma acc update self (a[0:N], b[0:N]) + + for (i = 0; i < N; i++) + { + if (a[i] != 5.0) + abort (); + + if (b[i] != 5.0) + abort (); + } + + if (!acc_is_present (&a[0], (N * sizeof (float)))) + abort (); + + if (!acc_is_present (&b[0], (N * sizeof (float)))) + abort (); + + for (i = 0; i < N; i++) + { + a[i] = 5.0; + b[i] = 1.0; + } + +#pragma acc update device (a[0:N], b[0:N]) + +#pragma acc parallel present (a[0:N], b[0:N]) + { + int ii; + + for (ii = 0; ii < N; ii++) + b[ii] = a[ii]; + } + +#pragma acc update host (a[0:N], b[0:N]) + + for (i = 0; i < N; i++) + { + if (a[i] != 5.0) + abort (); + + if (b[i] != 5.0) + abort (); + } + + if (!acc_is_present (&a[0], (N * sizeof (float)))) + abort (); + + if (!acc_is_present (&b[0], (N * sizeof (float)))) + abort (); + + for (i = 0; i < N; i++) + { + a[i] = 6.0; + b[i] = 0.0; + } + +#pragma acc update device (a[0:N], b[0:N]) + + for (i = 0; i < N; i++) + { + a[i] = 9.0; + } + +#pragma acc parallel present (a[0:N], b[0:N]) + { + int ii; + + for (ii = 0; ii < N; ii++) + b[ii] = a[ii]; + } + +#pragma acc update self (a[0:N], b[0:N]) + + for (i = 0; i < N; i++) + { + if (a[i] != 6.0) + abort (); + + if (b[i] != 6.0) + abort (); + } + + if (!acc_is_present (&a[0], (N * sizeof (float)))) + abort (); + + if (!acc_is_present (&b[0], (N * sizeof (float)))) + abort (); + + for (i = 0; i < N; i++) + { + a[i] = 7.0; + b[i] = 2.0; + } + +#pragma acc update device (a[0:N], b[0:N]) + + for (i = 0; i < N; i++) + { + a[i] = 9.0; + } + +#pragma acc parallel present (a[0:N], b[0:N]) + { + int ii; + + for (ii = 0; ii < N; ii++) + b[ii] = a[ii]; + } + +#pragma acc update self (a[0:N], b[0:N]) + + for (i = 0; i < N; i++) + { + if (a[i] != 7.0) + abort (); + + if (b[i] != 7.0) + abort (); + } + + for (i = 0; i < N; i++) + { + a[i] = 9.0; + } + +#pragma acc update device (a[0:N]) + +#pragma acc parallel present (a[0:N], b[0:N]) + { + int ii; + + for (ii = 0; ii < N; ii++) + b[ii] = a[ii]; + } + +#pragma acc update self (a[0:N], b[0:N]) + + for (i = 0; i < N; i++) + { + if (a[i] != 9.0) + abort (); + + if (b[i] != 9.0) + abort (); + } + + if (!acc_is_present (&a[0], (N * sizeof (float)))) + abort (); + + if (!acc_is_present (&b[0], (N * sizeof (float)))) + abort (); + + for (i = 0; i < N; i++) + { + a[i] = 5.0; + } + +#pragma acc update device (a[0:N]) + + for (i = 0; i < N; i++) + { + a[i] = 6.0; + } + +#pragma acc update device (a[0:N >> 1]) + +#pragma acc parallel present (a[0:N], b[0:N]) + { + int ii; + + for (ii = 0; ii < N; ii++) + b[ii] = a[ii]; + } + +#pragma acc update self (a[0:N], b[0:N]) + + for (i = 0; i < (N >> 1); i++) + { + if (a[i] != 6.0) + abort (); + + if (b[i] != 6.0) + abort (); + } + + for (i = (N >> 1); i < N; i++) + { + if (a[i] != 5.0) + abort (); + + if (b[i] != 5.0) + abort (); + } + + if (!acc_is_present (&a[0], (N * sizeof (float)))) + abort (); + + if (!acc_is_present (&b[0], (N * sizeof (float)))) + abort (); +} + +int +main (int argc, char **argv) +{ + int N = 8; + float *a, *b, *c; + float *d_a, *d_b, *d_c; + int i; + + a = (float *) malloc (N * sizeof (float)); + b = (float *) malloc (N * sizeof (float)); + c = (float *) malloc (N * sizeof (float)); + + d_a = (float *) acc_malloc (N * sizeof (float)); + d_b = (float *) acc_malloc (N * sizeof (float)); + d_c = (float *) acc_malloc (N * sizeof (float)); + + test (N, a, b, c, d_a, d_b, d_c, i); + + acc_free (d_a); + acc_free (d_b); + acc_free (d_c); + + free (a); + free (b); + free (c); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-c++/update-template.C b/libgomp/testsuite/libgomp.oacc-c++/update-template.C new file mode 100644 index 0000000..ae21e73 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c++/update-template.C @@ -0,0 +1,301 @@ +/* Copy of update-1.c with self exchanged with host for #pragma acc update. + This exercises templates. */ + +/* { dg-do run } */ +/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */ + +#include +#include +#include +#include +#include + +template +int +test (T1 &N, T2 *(&a), T2 *(&b), T2 *(&c), T2 *(&d_a), T2 *(&d_b), + T2 *(&d_c), T1 &i) +{ + + + for (i = 0; i < N; i++) + { + a[i] = 3.0; + b[i] = 0.0; + } + + acc_map_data (a, d_a, N * sizeof (float)); + acc_map_data (b, d_b, N * sizeof (float)); + acc_map_data (c, d_c, N * sizeof (float)); + +#pragma acc update device (a[0:N], b[0:N]) + +#pragma acc parallel present (a[0:N], b[0:N]) + { + int ii; + + for (ii = 0; ii < N; ii++) + b[ii] = a[ii]; + } + +#pragma acc update self (a[0:N], b[0:N]) + + for (i = 0; i < N; i++) + { + if (a[i] != 3.0) + abort (); + + if (b[i] != 3.0) + abort (); + } + + if (!acc_is_present (&a[0], (N * sizeof (float)))) + abort (); + + if (!acc_is_present (&b[0], (N * sizeof (float)))) + abort (); + + for (i = 0; i < N; i++) + { + a[i] = 5.0; + b[i] = 1.0; + } + +#pragma acc update device (a[0:N], b[0:N]) + +#pragma acc parallel present (a[0:N], b[0:N]) + { + int ii; + + for (ii = 0; ii < N; ii++) + b[ii] = a[ii]; + } + +#pragma acc update self (a[0:N], b[0:N]) + + for (i = 0; i < N; i++) + { + if (a[i] != 5.0) + abort (); + + if (b[i] != 5.0) + abort (); + } + + if (!acc_is_present (&a[0], (N * sizeof (float)))) + abort (); + + if (!acc_is_present (&b[0], (N * sizeof (float)))) + abort (); + + for (i = 0; i < N; i++) + { + a[i] = 5.0; + b[i] = 1.0; + } + +#pragma acc update device (a[0:N], b[0:N]) + +#pragma acc parallel present (a[0:N], b[0:N]) + { + int ii; + + for (ii = 0; ii < N; ii++) + b[ii] = a[ii]; + } + +#pragma acc update host (a[0:N], b[0:N]) + + for (i = 0; i < N; i++) + { + if (a[i] != 5.0) + abort (); + + if (b[i] != 5.0) + abort (); + } + + if (!acc_is_present (&a[0], (N * sizeof (float)))) + abort (); + + if (!acc_is_present (&b[0], (N * sizeof (float)))) + abort (); + + for (i = 0; i < N; i++) + { + a[i] = 6.0; + b[i] = 0.0; + } + +#pragma acc update device (a[0:N], b[0:N]) + + for (i = 0; i < N; i++) + { + a[i] = 9.0; + } + +#pragma acc parallel present (a[0:N], b[0:N]) + { + int ii; + + for (ii = 0; ii < N; ii++) + b[ii] = a[ii]; + } + +#pragma acc update self (a[0:N], b[0:N]) + + for (i = 0; i < N; i++) + { + if (a[i] != 6.0) + abort (); + + if (b[i] != 6.0) + abort (); + } + + if (!acc_is_present (&a[0], (N * sizeof (float)))) + abort (); + + if (!acc_is_present (&b[0], (N * sizeof (float)))) + abort (); + + for (i = 0; i < N; i++) + { + a[i] = 7.0; + b[i] = 2.0; + } + +#pragma acc update device (a[0:N], b[0:N]) + + for (i = 0; i < N; i++) + { + a[i] = 9.0; + } + +#pragma acc parallel present (a[0:N], b[0:N]) + { + int ii; + + for (ii = 0; ii < N; ii++) + b[ii] = a[ii]; + } + +#pragma acc update self (a[0:N], b[0:N]) + + for (i = 0; i < N; i++) + { + if (a[i] != 7.0) + abort (); + + if (b[i] != 7.0) + abort (); + } + + for (i = 0; i < N; i++) + { + a[i] = 9.0; + } + +#pragma acc update device (a[0:N]) + +#pragma acc parallel present (a[0:N], b[0:N]) + { + int ii; + + for (ii = 0; ii < N; ii++) + b[ii] = a[ii]; + } + +#pragma acc update self (a[0:N], b[0:N]) + + for (i = 0; i < N; i++) + { + if (a[i] != 9.0) + abort (); + + if (b[i] != 9.0) + abort (); + } + + if (!acc_is_present (&a[0], (N * sizeof (float)))) + abort (); + + if (!acc_is_present (&b[0], (N * sizeof (float)))) + abort (); + + for (i = 0; i < N; i++) + { + a[i] = 5.0; + } + +#pragma acc update device (a[0:N]) + + for (i = 0; i < N; i++) + { + a[i] = 6.0; + } + +#pragma acc update device (a[0:N >> 1]) + +#pragma acc parallel present (a[0:N], b[0:N]) + { + int ii; + + for (ii = 0; ii < N; ii++) + b[ii] = a[ii]; + } + +#pragma acc update self (a[0:N], b[0:N]) + + for (i = 0; i < (N >> 1); i++) + { + if (a[i] != 6.0) + abort (); + + if (b[i] != 6.0) + abort (); + } + + for (i = (N >> 1); i < N; i++) + { + if (a[i] != 5.0) + abort (); + + if (b[i] != 5.0) + abort (); + } + + if (!acc_is_present (&a[0], (N * sizeof (float)))) + abort (); + + if (!acc_is_present (&b[0], (N * sizeof (float)))) + abort (); +} + +int +main (int argc, char **argv) +{ + int N = 8; + float *a, *b, *c; + float *d_a, *d_b, *d_c; + int i; + + a = (float *) malloc (N * sizeof (float)); + b = (float *) malloc (N * sizeof (float)); + c = (float *) malloc (N * sizeof (float)); + + d_a = (float *) acc_malloc (N * sizeof (float)); + d_b = (float *) acc_malloc (N * sizeof (float)); + d_c = (float *) acc_malloc (N * sizeof (float)); + + test (N, a, b, c, d_a, d_b, d_c, i); + + acc_free (d_a); + acc_free (d_b); + acc_free (d_c); + + free (a); + free (b); + free (c); + + return 0; +} 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 22cef6d..67c382c 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/asyncwait-1.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/asyncwait-1.c @@ -1,3 +1,5 @@ +/* Test acc_set_cuda_stream integration with async and wait. */ + /* { dg-do run { target openacc_nvidia_accel_selected } } */ /* { dg-additional-options "-lcuda" } */ @@ -37,7 +39,7 @@ main (int argc, char **argv) #pragma acc data copy (a[0:N]) copy (b[0:N]) copyin (N) { -#pragma acc parallel async +#pragma acc parallel async present (a[0:N], b[0:N], N) { int ii; @@ -67,7 +69,7 @@ main (int argc, char **argv) #pragma acc data copy (a[0:N]) copy (b[0:N]) copyin (N) { -#pragma acc parallel async (1) +#pragma acc parallel async (1) present (a[0:N], b[0:N], N) { int ii; @@ -99,7 +101,7 @@ main (int argc, char **argv) #pragma acc data copy (a[0:N]) copy (b[0:N]) copy (c[0:N]) copy (d[0:N]) copyin (N) { -#pragma acc parallel async (1) +#pragma acc parallel async (1) present (a[0:N], b[0:N], N) { int ii; @@ -107,7 +109,7 @@ main (int argc, char **argv) b[ii] = (a[ii] * a[ii] * a[ii]) / a[ii]; } -#pragma acc parallel async (1) +#pragma acc parallel async (1) present (a[0:N], c[0:N], N) { int ii; @@ -116,7 +118,7 @@ main (int argc, char **argv) } -#pragma acc parallel async (1) +#pragma acc parallel async (1) present (a[0:N], d[0:N], N) { int ii; @@ -155,7 +157,7 @@ main (int argc, char **argv) #pragma acc data copy (a[0:N], b[0:N], c[0:N], d[0:N], e[0:N]) copyin (N) { -#pragma acc parallel async (1) +#pragma acc parallel async (1) present (a[0:N], b[0:N], N) { int ii; @@ -163,7 +165,7 @@ main (int argc, char **argv) b[ii] = (a[ii] * a[ii] * a[ii]) / a[ii]; } -#pragma acc parallel async (1) +#pragma acc parallel async (1) present (a[0:N], c[0:N], N) { int ii; @@ -171,7 +173,7 @@ main (int argc, char **argv) c[ii] = (a[ii] + a[ii] + a[ii] + a[ii]) / a[ii]; } -#pragma acc parallel async (1) +#pragma acc parallel async (1) present (a[0:N], d[0:N], N) { int ii; @@ -179,7 +181,8 @@ main (int argc, char **argv) d[ii] = ((a[ii] * a[ii] + a[ii]) / a[ii]) - a[ii]; } -#pragma acc parallel wait (1) async (1) +#pragma acc parallel wait (1) async (1) present (a[0:N], b[0:N], c[0:N]) \ + present (d[0:N], e[0:N], N) { int ii; @@ -228,7 +231,7 @@ main (int argc, char **argv) #pragma acc data copy (a[0:N], b[0:N]) copyin (N) { -#pragma acc parallel async (1) +#pragma acc parallel async (1) present (a[0:N], b[0:N], N) { int ii; @@ -257,10 +260,11 @@ main (int argc, char **argv) d[i] = 0.0; } -#pragma acc data copy (a[0:N]) copy (b[0:N]) copy (c[0:N]) copy (d[0:N]) copyin (N) +#pragma acc data copy (a[0:N]) copy (b[0:N]) copy (c[0:N]) copy (d[0:N]) \ + copyin (N) { -#pragma acc parallel async (1) +#pragma acc parallel async (1) present (a[0:N], b[0:N], N) { int ii; @@ -268,7 +272,7 @@ main (int argc, char **argv) b[ii] = (a[ii] * a[ii] * a[ii]) / a[ii]; } -#pragma acc parallel async (1) +#pragma acc parallel async (1) present (a[0:N], c[0:N], N) { int ii; @@ -276,7 +280,7 @@ main (int argc, char **argv) c[ii] = (a[ii] + a[ii] + a[ii] + a[ii]) / a[ii]; } -#pragma acc parallel async (1) +#pragma acc parallel async (1) present (a[0:N], d[0:N], N) { int ii; @@ -315,7 +319,7 @@ main (int argc, char **argv) #pragma acc data copy (a[0:N], b[0:N], c[0:N], d[0:N], e[0:N]) copyin (N) { -#pragma acc parallel async (1) +#pragma acc parallel async (1) present (a[0:N], b[0:N], N) { int ii; @@ -323,7 +327,7 @@ main (int argc, char **argv) b[ii] = (a[ii] * a[ii] * a[ii]) / a[ii]; } -#pragma acc parallel async (1) +#pragma acc parallel async (1) present (a[0:N], c[0:N], N) { int ii; @@ -331,7 +335,7 @@ main (int argc, char **argv) c[ii] = (a[ii] + a[ii] + a[ii] + a[ii]) / a[ii]; } -#pragma acc parallel async (1) +#pragma acc parallel async (1) present (a[0:N], d[0:N], N) { int ii; @@ -339,7 +343,8 @@ main (int argc, char **argv) d[ii] = ((a[ii] * a[ii] + a[ii]) / a[ii]) - a[ii]; } -#pragma acc parallel wait (1) async (1) +#pragma acc parallel wait (1) async (1) present (a[0:N], b[0:N], c[0:N]) \ + present (d[0:N], e[0:N], N) { int ii; @@ -381,7 +386,7 @@ main (int argc, char **argv) #pragma acc data copyin (a[0:N], b[0:N], c[0:N]) copyin (N) { -#pragma acc parallel async (1) +#pragma acc parallel async (1) present (a[0:N], b[0:N], N) { int ii; @@ -389,7 +394,7 @@ main (int argc, char **argv) b[ii] = (a[ii] * a[ii] * a[ii]) / a[ii]; } -#pragma acc parallel async (1) +#pragma acc parallel async (1) present (a[0:N], c[0:N], N) { int ii; @@ -426,7 +431,7 @@ main (int argc, char **argv) #pragma acc data copyin (a[0:N], b[0:N], c[0:N]) copyin (N) { -#pragma acc parallel async (1) +#pragma acc parallel async (1) present (a[0:N], b[0:N], N) { int ii; @@ -434,7 +439,7 @@ main (int argc, char **argv) b[ii] = (a[ii] * a[ii] * a[ii]) / a[ii]; } -#pragma acc parallel async (1) +#pragma acc parallel async (1) present (a[0:N], c[0:N], N) { int ii; @@ -462,5 +467,11 @@ main (int argc, char **argv) acc_shutdown (acc_device_nvidia); + free (a); + free (b); + free (c); + free (d); + free (e); + 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..bbdaabe 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 data" and "acc exit data". */ + /* { 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 async wait present (a[0:N], b[0:N], N) #pragma acc loop for (i = 0; i < N; i++) b[i] = a[i]; @@ -41,7 +43,7 @@ main (int argc, char **argv) if (b[i] != 3.0) abort (); } - + return 0; for (i = 0; i < N; i++) { a[i] = 2.0; @@ -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 async (1) present (a[0:N], b[0:N], N) #pragma acc loop for (i = 0; i < N; i++) b[i] = a[i]; @@ -74,24 +76,26 @@ main (int argc, char **argv) d[i] = 0.0; } -#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 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 async (1) wait (1) present (a[0:N], b[0:N], N) #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 async (2) wait (1) present (a[0:N], c[0:N], N) #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 async (3) wait (1) present (a[0:N], d[0:N], N) #pragma acc loop for (i = 0; i < N; i++) d[i] = ((a[i] * a[i] + a[i]) / a[i]) - a[i]; -#pragma acc exit data copyout (a[0:N]) copyout (b[0:N]) copyout (c[0:N]) copyout (d[0:N]) wait (1, 2, 3) async (1) +#pragma acc exit data copyout (a[0:N]) copyout (b[0:N]) copyout (c[0:N]) \ + copyout (d[0:N]) wait (1, 2, 3) async (1) #pragma acc wait (1) for (i = 0; i < N; i++) @@ -118,28 +122,30 @@ main (int argc, char **argv) e[i] = 0.0; } -#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 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 async (1) wait (1) present (a[0:N], b[0:N], N) 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 async (2) wait (1) present (a[0:N], c[0:N], N) 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 async (3) wait (1) present (a[0:N], d[0:N], N) 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 wait (1) async (4) present (a[0:N], b[0:N], c[0:N]) \ + present (d[0:N], e[0:N], N) 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) delete (N) #pragma acc wait (1) - for (i = 0; i < N; i++) { if (a[i] != 2.0) @@ -158,5 +164,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/data-3.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/data-3.c index 747109f..014eb85 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 data", "acc exit data" and "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 async wait present (a[0:N], b[0:N], N) #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 async (1) present (a[0:N], b[0:N], N) #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 async (1) wait (1,2) present (a[0:N], b[0:N], N) #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 async (2) wait (1,3) present (a[0:N], c[0:N], N) #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 async (3) wait (1,3) present (a[0:N], d[0:N], N) #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 async (1) wait (1) present (a[0:N], b[0:N], N) 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 async (2) wait (1) present (a[0:N], c[0:N], N) 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 async (3) wait (1) present (a[0:N], d[0:N], N) 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 wait (1,5) async (4) present (a[0:N], b[0:N], c[0:N]) \ + present (d[0:N], e[0:N], N) 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/if-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/if-1.c index 6aa3bb7..cb6e59c 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/if-1.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/if-1.c @@ -457,7 +457,7 @@ main(int argc, char **argv) #pragma acc data copyin(a[0:N]) copyout(b[0:N]) if(1) { -#pragma acc parallel present(a[0:N]) +#pragma acc parallel present(a[0:N], b[0:N]) { int ii; diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/nested-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/nested-1.c index ededf2b..86dd491 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/nested-1.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/nested-1.c @@ -1,3 +1,5 @@ +/* Test implicit data clauses inside data regions. */ + /* { dg-do run } */ /* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */ @@ -11,20 +13,16 @@ int main (int argc, char **argv) { int N = 8; - float *a, *b, *c, *d; + float a[N], b[N], c[N], *d; int i; - a = (float *) malloc (N * sizeof (float)); - b = (float *) malloc (N * sizeof (float)); - c = (float *) malloc (N * sizeof (float)); - for (i = 0; i < N; i++) { a[i] = 3.0; b[i] = 0.0; } -#pragma acc data copyin (a[0:N]) copyout (b[0:N]) +#pragma acc data copyin (a) copyout (b) { #pragma acc parallel { @@ -53,7 +51,7 @@ main (int argc, char **argv) b[i] = 1.0; } -#pragma acc data copyin (a[0:N]) copyout (b[0:N]) +#pragma acc data copyin (a) copyout (b) { #pragma acc parallel { @@ -89,7 +87,7 @@ main (int argc, char **argv) a[i] = 9.0; } -#pragma acc data present_or_copyin (a[0:N]) copyout (b[0:N]) +#pragma acc data present_or_copyin (a) copyout (b) { #pragma acc parallel { @@ -120,7 +118,7 @@ main (int argc, char **argv) b[i] = 0.0; } -#pragma acc data copyin (a[0:N]) present_or_copyout (b[0:N]) +#pragma acc data copyin (a) present_or_copyout (b) { #pragma acc parallel { @@ -151,7 +149,7 @@ main (int argc, char **argv) d = (float *) acc_copyin (&b[0], N * sizeof (float)); -#pragma acc data copyin (a[0:N]) present_or_copyout (b[0:N]) +#pragma acc data copyin (a) present_or_copyout (b) { #pragma acc parallel { @@ -188,7 +186,7 @@ main (int argc, char **argv) b[i] = 4.0; } -#pragma acc data copy (a[0:N]) copyout (b[0:N]) +#pragma acc data copy (a) copyout (b) { #pragma acc parallel { @@ -223,7 +221,7 @@ main (int argc, char **argv) b[i] = 7.0; } -#pragma acc data present_or_copy (a[0:N]) present_or_copy (b[0:N]) +#pragma acc data present_or_copy (a) present_or_copy (b) { #pragma acc parallel { @@ -261,7 +259,7 @@ main (int argc, char **argv) d = (float *) acc_copyin (&a[0], N * sizeof (float)); d = (float *) acc_copyin (&b[0], N * sizeof (float)); -#pragma acc data present_or_copy (a[0:N]) present_or_copy (b[0:N]) +#pragma acc data present_or_copy (a) present_or_copy (b) { #pragma acc parallel { @@ -305,7 +303,7 @@ main (int argc, char **argv) b[i] = 7.0; } -#pragma acc data copyin (a[0:N]) create (c[0:N]) copyout (b[0:N]) +#pragma acc data copyin (a) create (c) copyout (b) { #pragma acc parallel { @@ -343,7 +341,7 @@ main (int argc, char **argv) b[i] = 8.0; } -#pragma acc data copyin (a[0:N]) present_or_create (c[0:N]) copyout (b[0:N]) +#pragma acc data copyin (a) present_or_create (c) copyout (b) { #pragma acc parallel { @@ -384,7 +382,7 @@ main (int argc, char **argv) d = (float *) acc_malloc (N * sizeof (float)); acc_map_data (c, d, N * sizeof (float)); -#pragma acc data copyin (a[0:N]) present_or_create (c[0:N]) copyout (b[0:N]) +#pragma acc data copyin (a) present_or_create (c) copyout (b) { #pragma acc parallel { @@ -431,7 +429,7 @@ main (int argc, char **argv) d = (float *) acc_malloc (N * sizeof (float)); acc_map_data (c, d, N * sizeof (float)); -#pragma acc data copyin (a[0:N]) present (c[0:N]) copyout (b[0:N]) +#pragma acc data copyin (a) present (c) copyout (b) { #pragma acc parallel { @@ -488,7 +486,7 @@ main (int argc, char **argv) if (!acc_is_present (a, (N * sizeof (float)))) abort (); -#pragma acc data present (a[0:N]) present (c[0:N]) present (b[0:N]) +#pragma acc data present (a) present (c) present (b) { #pragma acc parallel { @@ -543,7 +541,7 @@ main (int argc, char **argv) d = (float *) acc_malloc (N * sizeof (float)); -#pragma acc parallel copyin (a[0:N]) deviceptr (d) copyout (b[0:N]) +#pragma acc parallel copyin (a) deviceptr (d) copyout (b) { int ii; @@ -584,7 +582,7 @@ main (int argc, char **argv) a[i] = 9.0; } -#pragma acc data pcopyin (a[0:N]) copyout (b[0:N]) +#pragma acc data pcopyin (a) copyout (b) { #pragma acc parallel { @@ -615,7 +613,7 @@ main (int argc, char **argv) b[i] = 0.0; } -#pragma acc data copyin (a[0:N]) pcopyout (b[0:N]) +#pragma acc data copyin (a) pcopyout (b) { #pragma acc parallel { @@ -644,7 +642,7 @@ main (int argc, char **argv) b[i] = 7.0; } -#pragma acc data copyin (a[0:N]) pcreate (c[0:N]) copyout (b[0:N]) +#pragma acc data copyin (a) pcreate (c) copyout (b) { #pragma acc parallel { diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/present-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/present-2.c index 41efa70..0c03bad 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/present-2.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/present-2.c @@ -1,3 +1,5 @@ +/* Test both explicit and implicitly present data. */ + /* { dg-do run } */ /* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */ @@ -8,11 +10,10 @@ int main (int argc, char **argv) { int N = 8; - float *a, *b; + float *a, b[N]; int i; a = (float *) malloc (N * sizeof (float)); - b = (float *) malloc (N * sizeof (float)); for (i = 0; i < N; i++) { @@ -20,7 +21,7 @@ main (int argc, char **argv) b[i] = 0.0; } -#pragma acc data copyin(a[0:N]) copyout(b[0:N]) +#pragma acc data copyin(a[0:N]) copy (b) { #pragma acc parallel present(a[0:N]) @@ -44,5 +45,7 @@ main (int argc, char **argv) abort (); } + free (a); + return 0; }