From patchwork Thu Mar 24 16:57:55 2016 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 8bit X-Patchwork-Submitter: Thomas Schwinge X-Patchwork-Id: 601727 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 3qWCLd2Cp5z9sBG for ; Fri, 25 Mar 2016 03:58:32 +1100 (AEDT) Authentication-Results: ozlabs.org; dkim=pass (1024-bit key; unprotected) header.d=gcc.gnu.org header.i=@gcc.gnu.org header.b=I6IJICc7; 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:from :to:cc:subject:in-reply-to:references:date:message-id :mime-version:content-type; q=dns; s=default; b=w8hvfkdkZHMI1nA0 uKh8ePjvwEQZsQeStn3T+HBWb49uGV+YELiziJI851EFt2Uno0SE9UzabejY8DK/ xtO/Z/JQ/op/h5LmpgXJzGBuN+p6tmeowv8zzGfErlYyZfKXRKqUcSFpnnjegTsz GU7CDcMKmdMo6JyZ8JBdsMn8EJs= DKIM-Signature: v=1; a=rsa-sha1; c=relaxed; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender:from :to:cc:subject:in-reply-to:references:date:message-id :mime-version:content-type; s=default; bh=U7gdaG379q8smdvR5RYz+L 7AzE4=; b=I6IJICc7n8M5A4EKYEtXuld4zEtmL6d3GPtucpjH3fTgpi1YIYwb2R sNtK2Ep8IwcZn1Ni4Qw8PiORpCeEZzMCChOw7i/MZVJ4O4nSTeN9ZDc9Z/iwnk2w /EgbRf/LiD6QgOkSN+KrkutQUjdeVruVhKUvzVqASsKqkHi2h75hM= Received: (qmail 55077 invoked by alias); 24 Mar 2016 16:58:23 -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 55051 invoked by uid 89); 24 Mar 2016 16:58:20 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-1.7 required=5.0 tests=AWL, BAYES_00, RCVD_IN_DNSWL_NONE, SPF_PASS autolearn=ham version=3.3.2 spammy=julian, Julian, jul, Jul X-HELO: relay1.mentorg.com Received: from relay1.mentorg.com (HELO relay1.mentorg.com) (192.94.38.131) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with (AES256-GCM-SHA384 encrypted) ESMTPS; Thu, 24 Mar 2016 16:58:10 +0000 Received: from svr-orw-fem-06.mgc.mentorg.com ([147.34.97.120]) by relay1.mentorg.com with esmtp id 1aj8aF-0003Wv-SS from Thomas_Schwinge@mentor.com for gcc-patches@gcc.gnu.org; Thu, 24 Mar 2016 09:58:07 -0700 Received: from tftp-cs (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; Thu, 24 Mar 2016 09:58:07 -0700 Received: by tftp-cs (Postfix, from userid 49978) id 1D1A4C2222; Thu, 24 Mar 2016 09:58:07 -0700 (PDT) From: Thomas Schwinge To: CC: Julian Brown , , Cesar Philippidis Subject: Re: [gomp4] Some additional OpenACC reduction tests In-Reply-To: <20150729182312.588bc7c1@octopus> References: <20150729182312.588bc7c1@octopus> User-Agent: Notmuch/0.9-125-g4686d11 (http://notmuchmail.org) Emacs/24.5.1 (i586-pc-linux-gnu) Date: Thu, 24 Mar 2016 17:57:55 +0100 Message-ID: <87egazrey4.fsf@kepler.schwinge.homeip.net> MIME-Version: 1.0 Hi! On Wed, 29 Jul 2015 18:23:12 +0100, Julian Brown wrote: > This is a set of 19 new tests for OpenACC reductions, covering several > ways of performing reductions over the parallel and loop directives > using gang or worker/vector level parallelism. > --- /dev/null > +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-2.c > @@ -0,0 +1,40 @@ > +#include > + > +/* Test of reduction on both parallel and loop directives (workers and vectors > + in gang-partitioned mode, int type with XOR). */ > + > +int > +main (int argc, char *argv[]) > +{ > + int i, j, arr[32768], res = 0, hres = 0; > + > + for (i = 0; i < 32768; i++) > + arr[i] = i; > + > + #pragma acc parallel num_gangs(32) num_workers(32) vector_length(32) \ > + reduction(^:res) > + { > + #pragma acc loop gang > + for (j = 0; j < 32; j++) > + { > + #pragma acc loop worker vector reduction(^:res) > + for (i = 0; i < 1024; i++) > + res ^= arr[j * 1024 + i]; > + > + #pragma acc loop worker vector reduction(^:res) > + for (i = 0; i < 1024; i++) > + res ^= arr[j * 1024 + (1023 - i)]; > + } > + } > + > + for (j = 0; j < 32; j++) > + for (i = 0; i < 1024; i++) > + { > + hres ^= arr[j * 1024 + i]; > + hres ^= arr[j * 1024 + (1023 - i)]; > + } > + > + assert (res == hres); > + > + return 0; > +} Given the interpretation of the OpenACC specification that the current implementation of OpenACC reductions in GCC is base upon (which we're currently re-visiting), it had been neccessary to add data clauses next to parallel constructs' reduction clauses -- but not for this test case. I now found why; it just happend to ;-) always pass, because apparently the two XOR loops' iterations just cancelled their values, so in the end, we'd always get an "unremarkable" result of zero for both res and hres. In gomp-4_0-branch r234461, I have now committed the following: commit 8fff8ae7117c21d6b4a701a63cdd4634950418d1 Author: tschwinge Date: Thu Mar 24 16:54:55 2016 +0000 Improve libgomp.oacc-c-c++-common/par-loop-comb-reduction-2.c libgomp/ * testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-2.c: Make failure observable. Add data clause next to parallel construct's reduction clause. git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gomp-4_0-branch@234461 138bc75d-0d04-0410-961f-82ee72b054a4 --- libgomp/ChangeLog.gomp | 6 ++++++ .../testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-2.c | 6 +++--- 2 files changed, 9 insertions(+), 3 deletions(-) Grüße Thomas diff --git libgomp/ChangeLog.gomp libgomp/ChangeLog.gomp index 53ae315..b10ae94 100644 --- libgomp/ChangeLog.gomp +++ libgomp/ChangeLog.gomp @@ -1,3 +1,9 @@ +2016-03-24 Thomas Schwinge + + * testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-2.c: + Make failure observable. Add data clause next to parallel + construct's reduction clause. + 2016-03-11 Cesar Philippidis * testsuite/libgomp.oacc-c-c++-common/vprop.c: New test. diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-2.c libgomp/testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-2.c index a7a75a9..5e4590f 100644 --- libgomp/testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-2.c +++ libgomp/testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-2.c @@ -12,14 +12,14 @@ main (int argc, char *argv[]) arr[i] = i; #pragma acc parallel num_gangs(32) num_workers(32) vector_length(32) \ - reduction(^:res) + reduction(^:res) copy(res) { #pragma acc loop gang for (j = 0; j < 32; j++) { #pragma acc loop worker vector reduction(^:res) for (i = 0; i < 1024; i++) - res ^= arr[j * 1024 + i]; + res ^= 3 * arr[j * 1024 + i]; #pragma acc loop worker vector reduction(^:res) for (i = 0; i < 1024; i++) @@ -30,7 +30,7 @@ main (int argc, char *argv[]) for (j = 0; j < 32; j++) for (i = 0; i < 1024; i++) { - hres ^= arr[j * 1024 + i]; + hres ^= 3 * arr[j * 1024 + i]; hres ^= arr[j * 1024 + (1023 - i)]; }