From patchwork Fri Aug 14 16:42:01 2015 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Cesar Philippidis X-Patchwork-Id: 507509 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 CCC301401DE for ; Sat, 15 Aug 2015 02:42:18 +1000 (AEST) Authentication-Results: ozlabs.org; dkim=pass (1024-bit key; unprotected) header.d=gcc.gnu.org header.i=@gcc.gnu.org header.b=FLRoS3Az; 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 :message-id:date:from:mime-version:to:subject:content-type; q= dns; s=default; b=mq5wA2OdZQO9G6Y3cqaQLqEV6an86lyls2lPJOuTf7o8GY CL7U5AAHdudHVOfWSxiGZMgTtRmMCkdYbbuYu8sH44vnaT0qt1VjN/jH786DHcFC eP7jPxlrAGCfwodbfuiwRDsJpl1g1x6gbuCmXGdNlSiYhAOIt9gn98HmWnIPs= 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 :message-id:date:from:mime-version:to:subject:content-type; s= default; bh=HtZICsgAEw76uTSLBVHTzXAPDfo=; b=FLRoS3Az9IlPIgLceWhB lYJJPgPLgPQ5qbjDdoKA0l4PSjBl7E1O7pwGcMjMVhMhYl15DE6dK4t0ot8WbLxu Ag9Fnrww3kHZfh8VvTFnpSaWChUDFOpoIvXEyJcQ37u/SwGNEnp7ru4/Nd5XzLT3 +XPqELC0tdSurAfhGXj71bY= Received: (qmail 43996 invoked by alias); 14 Aug 2015 16:42:12 -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 43981 invoked by uid 89); 14 Aug 2015 16:42:11 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-2.0 required=5.0 tests=AWL, BAYES_00, RCVD_IN_DNSWL_LOW, SPF_PASS autolearn=ham version=3.3.2 X-HELO: relay1.mentorg.com Received: from relay1.mentorg.com (HELO relay1.mentorg.com) (192.94.38.131) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with ESMTP; Fri, 14 Aug 2015 16:42:05 +0000 Received: from svr-orw-fem-06.mgc.mentorg.com ([147.34.97.120]) by relay1.mentorg.com with esmtp id 1ZQI3N-0000kv-UL from Cesar_Philippidis@mentor.com for gcc-patches@gcc.gnu.org; Fri, 14 Aug 2015 09:42:01 -0700 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; Fri, 14 Aug 2015 09:42:01 -0700 Message-ID: <55CE1A59.3070709@codesourcery.com> Date: Fri, 14 Aug 2015 09:42:01 -0700 From: Cesar Philippidis User-Agent: Mozilla/5.0 (X11; Linux x86_64; rv:31.0) Gecko/20100101 Thunderbird/31.8.0 MIME-Version: 1.0 To: "gcc-patches@gcc.gnu.org" Subject: [gomp4] non-acc loop reductions implicit copy bugfix This patch teaches the c and c++ front ends not to add a 'copy' clause for each non-loop OpenACC reduction. Without this patch, this construct would error with a duplicate mapping for 'sum': #pragma acc parallel num_gangs (10) copy (sum) reduction (+:sum) While the intent behind adding a 'copy' clause for the reduction variable was good, it is inconsistent with the behavior of the other data clauses in the OpenACC spec. Without an explicit copy, 'sum' should probably be transferred to the accelerator as firstprivate. Anyway, it's probably better to correct this behavior in the gimplifier later on if necessary. I'll apply this patch to gomp-4_0-branch shortly. Cesar 2015-08-14 Cesar Philippidis gcc/ * c/c-typeck.c (c_finish_omp_clauses): Permit variables to appear in both OpenACC data and reduction clauses. * cp/semantics.c (finish_omp_clauses): Likewise. gcc/testsuite/ * c-c++-common/goacc/parallel-reduction.c: New test. diff --git a/gcc/c/c-typeck.c b/gcc/c/c-typeck.c index ed748c8..fe54799f 100644 --- a/gcc/c/c-typeck.c +++ b/gcc/c/c-typeck.c @@ -12059,11 +12059,10 @@ tree c_finish_omp_clauses (tree clauses, bool oacc) { bitmap_head generic_head, firstprivate_head, lastprivate_head; - bitmap_head aligned_head, oacc_data_head; + bitmap_head aligned_head, oacc_data_head, oacc_reduction_head; tree c, t, *pc; bool branch_seen = false; bool copyprivate_seen = false; - bool oacc_data = false; tree *nowait_clause = NULL; bitmap_obstack_initialize (NULL); @@ -12072,12 +12071,15 @@ c_finish_omp_clauses (tree clauses, bool oacc) bitmap_initialize (&lastprivate_head, &bitmap_default_obstack); bitmap_initialize (&aligned_head, &bitmap_default_obstack); bitmap_initialize (&oacc_data_head, &bitmap_default_obstack); + bitmap_initialize (&oacc_reduction_head, &bitmap_default_obstack); for (pc = &clauses, c = clauses; c ; c = *pc) { bool remove = false; bool need_complete = false; bool need_implicitly_determined = false; + bool oacc_data = false; + bool reduction = false; switch (OMP_CLAUSE_CODE (c)) { @@ -12095,8 +12097,8 @@ c_finish_omp_clauses (tree clauses, bool oacc) goto check_dup_generic; case OMP_CLAUSE_REDUCTION: - need_implicitly_determined = true; - oacc_data = false; + need_implicitly_determined = !oacc; + reduction = true; t = OMP_CLAUSE_DECL (c); if (OMP_CLAUSE_REDUCTION_PLACEHOLDER (c) == NULL_TREE && (FLOAT_TYPE_P (TREE_TYPE (t)) @@ -12312,12 +12314,23 @@ c_finish_omp_clauses (tree clauses, bool oacc) else bitmap_set_bit (&oacc_data_head, DECL_UID (t)); } + else if (reduction) + { + if (oacc && bitmap_bit_p (&oacc_reduction_head, DECL_UID (t))) + { + error_at (OMP_CLAUSE_LOCATION (c), + "%qE appears in multiple 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))) { error_at (OMP_CLAUSE_LOCATION (c), - "%qE appears more than once in data clauses", t); + "%qE appears more than one non-data clause", t); remove = true; } else diff --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c index 0b13ca2..cf1790c 100644 --- a/gcc/cp/semantics.c +++ b/gcc/cp/semantics.c @@ -5281,11 +5281,10 @@ tree finish_omp_clauses (tree clauses, bool oacc) { bitmap_head generic_head, firstprivate_head, lastprivate_head; - bitmap_head aligned_head, oacc_data_head; + bitmap_head aligned_head, oacc_data_head, oacc_reduction_head; tree c, t, *pc; bool branch_seen = false; bool copyprivate_seen = false; - bool oacc_data = false; bitmap_obstack_initialize (NULL); bitmap_initialize (&generic_head, &bitmap_default_obstack); @@ -5293,10 +5292,13 @@ finish_omp_clauses (tree clauses, bool oacc) bitmap_initialize (&lastprivate_head, &bitmap_default_obstack); bitmap_initialize (&aligned_head, &bitmap_default_obstack); bitmap_initialize (&oacc_data_head, &bitmap_default_obstack); + bitmap_initialize (&oacc_reduction_head, &bitmap_default_obstack); for (pc = &clauses, c = clauses; c ; c = *pc) { bool remove = false; + bool oacc_data = false; + bool reduction = false; switch (OMP_CLAUSE_CODE (c)) { @@ -5314,6 +5316,7 @@ finish_omp_clauses (tree clauses, bool oacc) if (oacc) { oacc_data = false; + reduction = true; goto check_dup_oacc; } else @@ -5426,6 +5429,17 @@ finish_omp_clauses (tree clauses, bool oacc) else bitmap_set_bit (&oacc_data_head, DECL_UID (t)); } + else if (reduction) + { + if (oacc && bitmap_bit_p (&oacc_reduction_head, DECL_UID (t))) + { + error_at (OMP_CLAUSE_LOCATION (c), + "%qE appears in multiple 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))) 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..debed55 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/parallel-reduction.c @@ -0,0 +1,13 @@ +int +main () +{ + int sum = 0; + +#pragma acc parallel num_gangs (10) copy (sum) reduction (+:sum) + { + int v = 5; + sum += 10 + v; + } + + return sum; +}