From patchwork Tue Apr 12 11:08:09 2016 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 8bit X-Patchwork-Submitter: Thomas Schwinge X-Patchwork-Id: 609344 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 3qkkh90q1bz9s2G for ; Tue, 12 Apr 2016 21:08:40 +1000 (AEST) Authentication-Results: ozlabs.org; dkim=pass (1024-bit key; unprotected) header.d=gcc.gnu.org header.i=@gcc.gnu.org header.b=Ka6KSd+o; 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:content-transfer-encoding; q=dns; s= default; b=JYzRu9rzc62C1VJ+tWo6gBgXbZBwCWsHOSDSzDJA3gY7oT8DDjIgl jP9XQRB+L7x9KJNq0qcNQlRGaH85x1RIRh2JlC6NeHcknN9zLxdzNRULi93m+WU1 98x7/UUZpVXYAyO9Q+lJYEpJ91chG2nF0uXie8zyixDLUHqwDXgh4o= 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:content-transfer-encoding; s=default; bh=teuP409TY1mGkaNJVYYaYO+7Tr0=; b=Ka6KSd+od3BEcwyw1NneX+0ovUAt syonUNAs5FdboZaEwATB/k7cRcivc7uoHTWFHVZu/L62CxbXz3bzeAhnHGW33SQc /rbVvVb4LB0PQ/rPI5rAe4kvV+DoiWH6ig/ezzxhzvZ4iS0EhTguLJQ3+0N22ejb HJDa+qIHRvgXY7M= Received: (qmail 60423 invoked by alias); 12 Apr 2016 11:08:30 -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 60410 invoked by uid 89); 12 Apr 2016 11:08:29 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-1.9 required=5.0 tests=AWL, BAYES_00, RCVD_IN_DNSWL_NONE, SPF_PASS autolearn=ham version=3.3.2 spammy=1, 55, inb, his, ina X-HELO: relay1.mentorg.com Received: from relay1.mentorg.com (HELO relay1.mentorg.com) (192.94.38.131) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with (AES256-GCM-SHA384 encrypted) ESMTPS; Tue, 12 Apr 2016 11:08:19 +0000 Received: from nat-ies.mentorg.com ([192.94.31.2] helo=SVR-IES-FEM-01.mgc.mentorg.com) by relay1.mentorg.com with esmtp id 1apwB6-0002qP-H5 from Thomas_Schwinge@mentor.com for gcc-patches@gcc.gnu.org; Tue, 12 Apr 2016 04:08:17 -0700 Received: from hertz.schwinge.homeip.net (137.202.0.76) by SVR-IES-FEM-01.mgc.mentorg.com (137.202.0.104) with Microsoft SMTP Server id 14.3.224.2; Tue, 12 Apr 2016 12:08:15 +0100 From: Thomas Schwinge To: CC: Cesar Philippidis Subject: Merge libgomp.oacc-c-c++-common/loop-reduction-*.c into libgomp.oacc-c-c++-common/reduction-7.c (was: [gomp4] Update OpenACC test cases) In-Reply-To: <87wpod8xp2.fsf@hertz.schwinge.homeip.net> References: <87bn5w9i1l.fsf@hertz.schwinge.homeip.net> <20160330141332.GD3017@tucnak.redhat.com> <878u109ew4.fsf@hertz.schwinge.homeip.net> <87wpod8xp2.fsf@hertz.schwinge.homeip.net> User-Agent: Notmuch/0.9-101-g81dad07 (http://notmuchmail.org) Emacs/24.4.1 (x86_64-pc-linux-gnu) Date: Tue, 12 Apr 2016 13:08:09 +0200 Message-ID: <87wpo33x0m.fsf@hertz.schwinge.homeip.net> MIME-Version: 1.0 Hi! On Mon, 04 Apr 2016 12:39:37 +0200, I wrote: > [...] gomp-4_0-branch [...] additional (cleanup) changes [...] > libgomp/ > [...] > * testsuite/libgomp.oacc-c-c++-common/loop-reduction-gang-np-1.c: > Merge this file, and... > * testsuite/libgomp.oacc-c-c++-common/loop-reduction-gv-np-1.c: > ... this file, and... > * testsuite/libgomp.oacc-c-c++-common/loop-reduction-gw-np-1.c: > ... this file, and... > * testsuite/libgomp.oacc-c-c++-common/loop-reduction-gwv-np-1.c: > ... this file, and... > * testsuite/libgomp.oacc-c-c++-common/loop-reduction-gwv-np-2.c: > ... this file, and... > * testsuite/libgomp.oacc-c-c++-common/loop-reduction-gwv-np-3.c: > ... this file, and... > * testsuite/libgomp.oacc-c-c++-common/loop-reduction-gwv-np-4.c: > ... this file, and... > * testsuite/libgomp.oacc-c-c++-common/loop-reduction-vector-p-1.c: > ... this file, and... > * testsuite/libgomp.oacc-c-c++-common/loop-reduction-vector-p-2.c: > ... this file, and... > * testsuite/libgomp.oacc-c-c++-common/loop-reduction-worker-p-1.c: > ... this file, and... > * testsuite/libgomp.oacc-c-c++-common/loop-reduction-wv-p-1.c: > ... this file, and... > * testsuite/libgomp.oacc-c-c++-common/loop-reduction-wv-p-2.c: > ... this file, and... > * testsuite/libgomp.oacc-c-c++-common/loop-reduction-wv-p-3.c: > ... this file into... > * testsuite/libgomp.oacc-c-c++-common/reduction-7.c: ... this new > file. Use dg-warning directives instead of specifying the -w > compiler option. > [...] Cesar didn't pick these up in his recent trunk commit, so now applied this to trunk in r234899: commit 40495bd0847a05aa76cc37e05292cf937449f9dd Author: tschwinge Date: Tue Apr 12 11:02:32 2016 +0000 Merge libgomp.oacc-c-c++-common/loop-reduction-*.c into libgomp.oacc-c-c++-common/reduction-7.c libgomp/ * testsuite/libgomp.oacc-c-c++-common/loop-reduction-gang-np-1.c: Merge this file, and... * testsuite/libgomp.oacc-c-c++-common/loop-reduction-gv-np-1.c: ... this file, and... * testsuite/libgomp.oacc-c-c++-common/loop-reduction-gw-np-1.c: ... this file, and... * testsuite/libgomp.oacc-c-c++-common/loop-reduction-gwv-np-1.c: ... this file, and... * testsuite/libgomp.oacc-c-c++-common/loop-reduction-gwv-np-2.c: ... this file, and... * testsuite/libgomp.oacc-c-c++-common/loop-reduction-gwv-np-3.c: ... this file, and... * testsuite/libgomp.oacc-c-c++-common/loop-reduction-gwv-np-4.c: ... this file, and... * testsuite/libgomp.oacc-c-c++-common/loop-reduction-vector-p-1.c: ... this file, and... * testsuite/libgomp.oacc-c-c++-common/loop-reduction-vector-p-2.c: ... this file, and... * testsuite/libgomp.oacc-c-c++-common/loop-reduction-worker-p-1.c: ... this file, and... * testsuite/libgomp.oacc-c-c++-common/loop-reduction-wv-p-1.c: ... this file, and... * testsuite/libgomp.oacc-c-c++-common/loop-reduction-wv-p-2.c: ... this file, and... * testsuite/libgomp.oacc-c-c++-common/loop-reduction-wv-p-3.c: ... this file into... * testsuite/libgomp.oacc-c-c++-common/reduction-7.c: ... this file. git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/trunk@234899 138bc75d-0d04-0410-961f-82ee72b054a4 --- libgomp/ChangeLog | 29 ++ .../loop-reduction-gang-np-1.c | 45 --- .../loop-reduction-gv-np-1.c | 30 -- .../loop-reduction-gw-np-1.c | 30 -- .../loop-reduction-gwv-np-1.c | 28 -- .../loop-reduction-gwv-np-2.c | 34 -- .../loop-reduction-gwv-np-3.c | 33 -- .../loop-reduction-gwv-np-4.c | 55 ---- .../loop-reduction-vector-p-1.c | 43 --- .../loop-reduction-vector-p-2.c | 41 --- .../loop-reduction-worker-p-1.c | 43 --- .../loop-reduction-wv-p-1.c | 41 --- .../loop-reduction-wv-p-2.c | 45 --- .../loop-reduction-wv-p-3.c | 38 --- .../libgomp.oacc-c-c++-common/reduction-7.c | 351 +++++++++++++++++++++ 15 files changed, 380 insertions(+), 506 deletions(-) Grüße Thomas diff --git libgomp/ChangeLog libgomp/ChangeLog index 6071b23..1716ba0 100644 --- libgomp/ChangeLog +++ libgomp/ChangeLog @@ -1,5 +1,34 @@ 2016-04-12 Thomas Schwinge + * testsuite/libgomp.oacc-c-c++-common/loop-reduction-gang-np-1.c: + Merge this file, and... + * testsuite/libgomp.oacc-c-c++-common/loop-reduction-gv-np-1.c: + ... this file, and... + * testsuite/libgomp.oacc-c-c++-common/loop-reduction-gw-np-1.c: + ... this file, and... + * testsuite/libgomp.oacc-c-c++-common/loop-reduction-gwv-np-1.c: + ... this file, and... + * testsuite/libgomp.oacc-c-c++-common/loop-reduction-gwv-np-2.c: + ... this file, and... + * testsuite/libgomp.oacc-c-c++-common/loop-reduction-gwv-np-3.c: + ... this file, and... + * testsuite/libgomp.oacc-c-c++-common/loop-reduction-gwv-np-4.c: + ... this file, and... + * testsuite/libgomp.oacc-c-c++-common/loop-reduction-vector-p-1.c: + ... this file, and... + * testsuite/libgomp.oacc-c-c++-common/loop-reduction-vector-p-2.c: + ... this file, and... + * testsuite/libgomp.oacc-c-c++-common/loop-reduction-worker-p-1.c: + ... this file, and... + * testsuite/libgomp.oacc-c-c++-common/loop-reduction-wv-p-1.c: + ... this file, and... + * testsuite/libgomp.oacc-c-c++-common/loop-reduction-wv-p-2.c: + ... this file, and... + * testsuite/libgomp.oacc-c-c++-common/loop-reduction-wv-p-3.c: + ... this file into... + * testsuite/libgomp.oacc-c-c++-common/reduction-7.c: ... this + file. + * testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-2.c: Make failure observable. diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/loop-reduction-gang-np-1.c libgomp/testsuite/libgomp.oacc-c-c++-common/loop-reduction-gang-np-1.c deleted file mode 100644 index 55ab3c9..0000000 --- libgomp/testsuite/libgomp.oacc-c-c++-common/loop-reduction-gang-np-1.c +++ /dev/null @@ -1,45 +0,0 @@ -/* { dg-additional-options "-w" } */ - -#include - -/* Test of reduction on loop directive (gangs, non-private reduction - variable). */ - -int -main (int argc, char *argv[]) -{ - int i, arr[1024], res = 0, hres = 0; - - for (i = 0; i < 1024; i++) - arr[i] = i; - - #pragma acc parallel num_gangs(32) num_workers(32) vector_length(32) \ - copy(res) - { - #pragma acc loop gang reduction(+:res) - for (i = 0; i < 1024; i++) - res += arr[i]; - } - - for (i = 0; i < 1024; i++) - hres += arr[i]; - - assert (res == hres); - - res = hres = 1; - - #pragma acc parallel num_gangs(32) num_workers(32) vector_length(32) \ - copy(res) - { - #pragma acc loop gang reduction(*:res) - for (i = 0; i < 12; i++) - res *= arr[i]; - } - - for (i = 0; i < 12; i++) - hres *= arr[i]; - - assert (res == hres); - - return 0; -} diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/loop-reduction-gv-np-1.c libgomp/testsuite/libgomp.oacc-c-c++-common/loop-reduction-gv-np-1.c deleted file mode 100644 index d4341e9..0000000 --- libgomp/testsuite/libgomp.oacc-c-c++-common/loop-reduction-gv-np-1.c +++ /dev/null @@ -1,30 +0,0 @@ -/* { dg-additional-options "-w" } */ - -#include - -/* Test of reduction on loop directive (gangs and vectors, non-private - reduction variable). */ - -int -main (int argc, char *argv[]) -{ - int i, arr[1024], res = 0, hres = 0; - - for (i = 0; i < 1024; i++) - arr[i] = i; - - #pragma acc parallel num_gangs(32) num_workers(32) vector_length(32) \ - copy(res) - { - #pragma acc loop gang vector reduction(+:res) - for (i = 0; i < 1024; i++) - res += arr[i]; - } - - for (i = 0; i < 1024; i++) - hres += arr[i]; - - assert (res == hres); - - return 0; -} diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/loop-reduction-gw-np-1.c libgomp/testsuite/libgomp.oacc-c-c++-common/loop-reduction-gw-np-1.c deleted file mode 100644 index 2e5668b..0000000 --- libgomp/testsuite/libgomp.oacc-c-c++-common/loop-reduction-gw-np-1.c +++ /dev/null @@ -1,30 +0,0 @@ -/* { dg-additional-options "-w" } */ - -#include - -/* Test of reduction on loop directive (gangs and workers, non-private - reduction variable). */ - -int -main (int argc, char *argv[]) -{ - int i, arr[1024], res = 0, hres = 0; - - for (i = 0; i < 1024; i++) - arr[i] = i; - - #pragma acc parallel num_gangs(32) num_workers(32) vector_length(32) \ - copy(res) - { - #pragma acc loop gang worker reduction(+:res) - for (i = 0; i < 1024; i++) - res += arr[i]; - } - - for (i = 0; i < 1024; i++) - hres += arr[i]; - - assert (res == hres); - - return 0; -} diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/loop-reduction-gwv-np-1.c libgomp/testsuite/libgomp.oacc-c-c++-common/loop-reduction-gwv-np-1.c deleted file mode 100644 index d610373..0000000 --- libgomp/testsuite/libgomp.oacc-c-c++-common/loop-reduction-gwv-np-1.c +++ /dev/null @@ -1,28 +0,0 @@ -#include - -/* Test of reduction on loop directive (gangs, workers and vectors, non-private - reduction variable). */ - -int -main (int argc, char *argv[]) -{ - int i, arr[1024], res = 0, hres = 0; - - for (i = 0; i < 1024; i++) - arr[i] = i; - - #pragma acc parallel num_gangs(32) num_workers(32) vector_length(32) \ - copy(res) - { - #pragma acc loop gang worker vector reduction(+:res) - for (i = 0; i < 1024; i++) - res += arr[i]; - } - - for (i = 0; i < 1024; i++) - hres += arr[i]; - - assert (res == hres); - - return 0; -} diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/loop-reduction-gwv-np-2.c libgomp/testsuite/libgomp.oacc-c-c++-common/loop-reduction-gwv-np-2.c deleted file mode 100644 index ea5c151..0000000 --- libgomp/testsuite/libgomp.oacc-c-c++-common/loop-reduction-gwv-np-2.c +++ /dev/null @@ -1,34 +0,0 @@ -#include - -/* Test of reduction on loop directive (gangs, workers and vectors, non-private - reduction variable: separate gang and worker/vector loops). */ - -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) \ - copy(res) - { - #pragma acc loop gang reduction(+:res) - 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" is non-private, and is not available until after the parallel - region. */ - } - - for (i = 0; i < 32768; i++) - hres += arr[i]; - - assert (res == hres); - - return 0; -} diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/loop-reduction-gwv-np-3.c libgomp/testsuite/libgomp.oacc-c-c++-common/loop-reduction-gwv-np-3.c deleted file mode 100644 index 0056f3c..0000000 --- libgomp/testsuite/libgomp.oacc-c-c++-common/loop-reduction-gwv-np-3.c +++ /dev/null @@ -1,33 +0,0 @@ -#include - -/* Test of reduction on loop directive (gangs, workers and vectors, non-private - reduction variable: separate gang and worker/vector loops). */ - -int -main (int argc, char *argv[]) -{ - int i, j; - double 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) \ - copyin(arr) copy(res) - { - #pragma acc loop gang reduction(+:res) - for (j = 0; j < 32; j++) - { - #pragma acc loop worker vector reduction(+:res) - for (i = 0; i < 1024; i++) - res += arr[j * 1024 + i]; - } - } - - for (i = 0; i < 32768; i++) - hres += arr[i]; - - assert (res == hres); - - return 0; -} diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/loop-reduction-gwv-np-4.c libgomp/testsuite/libgomp.oacc-c-c++-common/loop-reduction-gwv-np-4.c deleted file mode 100644 index e69d0ec..0000000 --- libgomp/testsuite/libgomp.oacc-c-c++-common/loop-reduction-gwv-np-4.c +++ /dev/null @@ -1,55 +0,0 @@ -#include - -/* Test of reduction on loop directive (gangs, workers and vectors, multiple - non-private reduction variables, float type). */ - -int -main (int argc, char *argv[]) -{ - int i, j; - float arr[32768]; - float res = 0, mres = 0, hres = 0, hmres = 0; - - for (i = 0; i < 32768; i++) - arr[i] = i; - - #pragma acc parallel num_gangs(32) num_workers(32) vector_length(32) \ - copy(res, mres) - { - #pragma acc loop gang reduction(+:res) reduction(max:mres) - for (j = 0; j < 32; j++) - { - #pragma acc loop worker vector reduction(+:res) reduction(max:mres) - for (i = 0; i < 1024; i++) - { - res += arr[j * 1024 + i]; - if (arr[j * 1024 + i] > mres) - mres = arr[j * 1024 + i]; - } - - #pragma acc loop worker vector reduction(+:res) reduction(max:mres) - for (i = 0; i < 1024; i++) - { - res += arr[j * 1024 + (1023 - i)]; - if (arr[j * 1024 + (1023 - i)] > mres) - mres = 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)]; - if (arr[j * 1024 + i] > hmres) - hmres = arr[j * 1024 + i]; - if (arr[j * 1024 + (1023 - i)] > hmres) - hmres = arr[j * 1024 + (1023 - i)]; - } - - assert (res == hres); - assert (mres == hmres); - - return 0; -} diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/loop-reduction-vector-p-1.c libgomp/testsuite/libgomp.oacc-c-c++-common/loop-reduction-vector-p-1.c deleted file mode 100644 index 31e4366..0000000 --- libgomp/testsuite/libgomp.oacc-c-c++-common/loop-reduction-vector-p-1.c +++ /dev/null @@ -1,43 +0,0 @@ -/* { dg-additional-options "-w" } */ - -#include - -/* Test of reduction on loop directive (vectors, private reduction - variable). */ - -int -main (int argc, char *argv[]) -{ - int i, j, arr[1024], out[32], res = 0, hres = 0; - - for (i = 0; i < 1024; i++) - arr[i] = i; - - #pragma acc parallel num_gangs(32) num_workers(32) vector_length(32) \ - private(res) copyout(out) - { - #pragma acc loop gang - for (j = 0; j < 32; j++) - { - res = 0; - - #pragma acc loop vector reduction(+:res) - for (i = 0; i < 32; i++) - res += arr[j * 32 + i]; - - out[j] = res; - } - } - - for (j = 0; j < 32; j++) - { - hres = 0; - - for (i = 0; i < 32; i++) - hres += arr[j * 32 + i]; - - assert (out[j] == hres); - } - - return 0; -} diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/loop-reduction-vector-p-2.c libgomp/testsuite/libgomp.oacc-c-c++-common/loop-reduction-vector-p-2.c deleted file mode 100644 index 15f0053..0000000 --- libgomp/testsuite/libgomp.oacc-c-c++-common/loop-reduction-vector-p-2.c +++ /dev/null @@ -1,41 +0,0 @@ -#include - -/* Test of reduction on loop directive (vector reduction in - gang-partitioned/worker-partitioned mode, private reduction variable). */ - -int -main (int argc, char *argv[]) -{ - int i, j, k; - double ina[1024], inb[1024], out[1024], acc; - - for (j = 0; j < 32; j++) - for (i = 0; i < 32; i++) - { - ina[j * 32 + i] = (i == j) ? 2.0 : 0.0; - inb[j * 32 + i] = (double) (i + j); - } - - #pragma acc parallel num_gangs(32) num_workers(32) vector_length(32) \ - private(acc) copyin(ina, inb) copyout(out) - { - #pragma acc loop gang worker - for (k = 0; k < 32; k++) - for (j = 0; j < 32; j++) - { - acc = 0; - - #pragma acc loop vector reduction(+:acc) - for (i = 0; i < 32; i++) - acc += ina[k * 32 + i] * inb[i * 32 + j]; - - out[k * 32 + j] = acc; - } - } - - for (j = 0; j < 32; j++) - for (i = 0; i < 32; i++) - assert (out[j * 32 + i] == (i + j) * 2); - - return 0; -} diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/loop-reduction-worker-p-1.c libgomp/testsuite/libgomp.oacc-c-c++-common/loop-reduction-worker-p-1.c deleted file mode 100644 index 4a92503..0000000 --- libgomp/testsuite/libgomp.oacc-c-c++-common/loop-reduction-worker-p-1.c +++ /dev/null @@ -1,43 +0,0 @@ -/* { dg-additional-options "-w" } */ - -#include - -/* Test of reduction on loop directive (workers, private reduction - variable). */ - -int -main (int argc, char *argv[]) -{ - int i, j, arr[1024], out[32], res = 0, hres = 0; - - for (i = 0; i < 1024; i++) - arr[i] = i; - - #pragma acc parallel num_gangs(32) num_workers(32) vector_length(32) \ - private(res) copyout(out) - { - #pragma acc loop gang - for (j = 0; j < 32; j++) - { - res = 0; - - #pragma acc loop worker reduction(+:res) - for (i = 0; i < 32; i++) - res += arr[j * 32 + i]; - - out[j] = res; - } - } - - for (j = 0; j < 32; j++) - { - hres = 0; - - for (i = 0; i < 32; i++) - hres += arr[j * 32 + i]; - - assert (out[j] == hres); - } - - return 0; -} diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/loop-reduction-wv-p-1.c libgomp/testsuite/libgomp.oacc-c-c++-common/loop-reduction-wv-p-1.c deleted file mode 100644 index 1bfb284..0000000 --- libgomp/testsuite/libgomp.oacc-c-c++-common/loop-reduction-wv-p-1.c +++ /dev/null @@ -1,41 +0,0 @@ -#include - -/* Test of reduction on loop directive (workers and vectors, private reduction - variable). */ - -int -main (int argc, char *argv[]) -{ - int i, j, arr[1024], out[32], res = 0, hres = 0; - - for (i = 0; i < 1024; i++) - arr[i] = i; - - #pragma acc parallel num_gangs(32) num_workers(32) vector_length(32) \ - private(res) copyout(out) - { - #pragma acc loop gang - for (j = 0; j < 32; j++) - { - res = 0; - - #pragma acc loop worker vector reduction(+:res) - for (i = 0; i < 32; i++) - res += arr[j * 32 + i]; - - out[j] = res; - } - } - - for (j = 0; j < 32; j++) - { - hres = 0; - - for (i = 0; i < 32; i++) - hres += arr[j * 32 + i]; - - assert (out[j] == hres); - } - - return 0; -} diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/loop-reduction-wv-p-2.c libgomp/testsuite/libgomp.oacc-c-c++-common/loop-reduction-wv-p-2.c deleted file mode 100644 index 93ab78f..0000000 --- libgomp/testsuite/libgomp.oacc-c-c++-common/loop-reduction-wv-p-2.c +++ /dev/null @@ -1,45 +0,0 @@ -#include - -/* Test of reduction on loop directive (workers and vectors, private reduction - variable). */ - -int -main (int argc, char *argv[]) -{ - int i, j, arr[32768], out[32], 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) \ - private(res) copyout(out) - { - #pragma acc loop gang - for (j = 0; j < 32; j++) - { - res = j; - - #pragma acc loop worker reduction(+:res) - for (i = 0; i < 1024; i++) - res += arr[j * 1024 + i]; - - #pragma acc loop vector reduction(+:res) - for (i = 1023; i >= 0; i--) - res += arr[j * 1024 + i]; - - out[j] = res; - } - } - - for (j = 0; j < 32; j++) - { - hres = j; - - for (i = 0; i < 1024; i++) - hres += arr[j * 1024 + i] * 2; - - assert (out[j] == hres); - } - - return 0; -} diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/loop-reduction-wv-p-3.c libgomp/testsuite/libgomp.oacc-c-c++-common/loop-reduction-wv-p-3.c deleted file mode 100644 index 298e25c..0000000 --- libgomp/testsuite/libgomp.oacc-c-c++-common/loop-reduction-wv-p-3.c +++ /dev/null @@ -1,38 +0,0 @@ -#include - -/* Test of reduction on loop directive (workers and vectors, private reduction - variable: gang-redundant mode). */ - -int -main (int argc, char *argv[]) -{ - int i, arr[1024], out[32], res = 0, hres = 0; - - for (i = 0; i < 1024; i++) - arr[i] = i ^ 33; - - #pragma acc parallel num_gangs(32) num_workers(32) vector_length(32) \ - private(res) copyin(arr) copyout(out) - { - /* Private variables aren't initialized by default in openacc. */ - res = 0; - - /* "res" should be available at the end of the following loop (and should - have the same value redundantly in each gang). */ - #pragma acc loop worker vector reduction(+:res) - for (i = 0; i < 1024; i++) - res += arr[i]; - - #pragma acc loop gang (static: 1) - for (i = 0; i < 32; i++) - out[i] = res; - } - - for (i = 0; i < 1024; i++) - hres += arr[i]; - - for (i = 0; i < 32; i++) - assert (out[i] == hres); - - return 0; -} diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/reduction-7.c libgomp/testsuite/libgomp.oacc-c-c++-common/reduction-7.c index b23c758..76c33e4 100644 --- libgomp/testsuite/libgomp.oacc-c-c++-common/reduction-7.c +++ libgomp/testsuite/libgomp.oacc-c-c++-common/reduction-7.c @@ -118,12 +118,363 @@ void gwv_np_1() } +/* Test of reduction on loop directive (gangs, workers and vectors, non-private + reduction variable: separate gang and worker/vector loops). */ + +void gwv_np_2() +{ + 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) \ + copy(res) + { + #pragma acc loop gang reduction(+:res) + 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" is non-private, and is not available until after the parallel + region. */ + } + + for (i = 0; i < 32768; i++) + hres += arr[i]; + + assert (res == hres); +} + + +/* Test of reduction on loop directive (gangs, workers and vectors, non-private + reduction variable: separate gang and worker/vector loops). */ + +void gwv_np_3() +{ + int i, j; + double 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) \ + copyin(arr) copy(res) + { + #pragma acc loop gang reduction(+:res) + for (j = 0; j < 32; j++) + { + #pragma acc loop worker vector reduction(+:res) + for (i = 0; i < 1024; i++) + res += arr[j * 1024 + i]; + } + } + + for (i = 0; i < 32768; i++) + hres += arr[i]; + + assert (res == hres); +} + + +/* Test of reduction on loop directive (gangs, workers and vectors, multiple + non-private reduction variables, float type). */ + +void gwv_np_4() +{ + int i, j; + float arr[32768]; + float res = 0, mres = 0, hres = 0, hmres = 0; + + for (i = 0; i < 32768; i++) + arr[i] = i; + + #pragma acc parallel num_gangs(32) num_workers(32) vector_length(32) \ + copy(res, mres) + { + #pragma acc loop gang reduction(+:res) reduction(max:mres) + for (j = 0; j < 32; j++) + { + #pragma acc loop worker vector reduction(+:res) reduction(max:mres) + for (i = 0; i < 1024; i++) + { + res += arr[j * 1024 + i]; + if (arr[j * 1024 + i] > mres) + mres = arr[j * 1024 + i]; + } + + #pragma acc loop worker vector reduction(+:res) reduction(max:mres) + for (i = 0; i < 1024; i++) + { + res += arr[j * 1024 + (1023 - i)]; + if (arr[j * 1024 + (1023 - i)] > mres) + mres = 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)]; + if (arr[j * 1024 + i] > hmres) + hmres = arr[j * 1024 + i]; + if (arr[j * 1024 + (1023 - i)] > hmres) + hmres = arr[j * 1024 + (1023 - i)]; + } + + assert (res == hres); + assert (mres == hmres); +} + + +/* Test of reduction on loop directive (vectors, private reduction + variable). */ + +void v_p_1() +{ + int i, j, arr[1024], out[32], res = 0, hres = 0; + + for (i = 0; i < 1024; i++) + arr[i] = i; + + #pragma acc parallel num_gangs(32) num_workers(32) vector_length(32) \ + private(res) copyout(out) + { + #pragma acc loop gang + for (j = 0; j < 32; j++) + { + res = 0; + + #pragma acc loop vector reduction(+:res) + for (i = 0; i < 32; i++) + res += arr[j * 32 + i]; + + out[j] = res; + } + } + + for (j = 0; j < 32; j++) + { + hres = 0; + + for (i = 0; i < 32; i++) + hres += arr[j * 32 + i]; + + assert (out[j] == hres); + } +} + + +/* Test of reduction on loop directive (vector reduction in + gang-partitioned/worker-partitioned mode, private reduction variable). */ + +void v_p_2() +{ + int i, j, k; + double ina[1024], inb[1024], out[1024], acc; + + for (j = 0; j < 32; j++) + for (i = 0; i < 32; i++) + { + ina[j * 32 + i] = (i == j) ? 2.0 : 0.0; + inb[j * 32 + i] = (double) (i + j); + } + + #pragma acc parallel num_gangs(32) num_workers(32) vector_length(32) \ + private(acc) copyin(ina, inb) copyout(out) + { + #pragma acc loop gang worker + for (k = 0; k < 32; k++) + for (j = 0; j < 32; j++) + { + acc = 0; + + #pragma acc loop vector reduction(+:acc) + for (i = 0; i < 32; i++) + acc += ina[k * 32 + i] * inb[i * 32 + j]; + + out[k * 32 + j] = acc; + } + } + + for (j = 0; j < 32; j++) + for (i = 0; i < 32; i++) + assert (out[j * 32 + i] == (i + j) * 2); +} + + +/* Test of reduction on loop directive (workers, private reduction + variable). */ + +void w_p_1() +{ + int i, j, arr[1024], out[32], res = 0, hres = 0; + + for (i = 0; i < 1024; i++) + arr[i] = i; + + #pragma acc parallel num_gangs(32) num_workers(32) vector_length(32) \ + private(res) copyout(out) + { + #pragma acc loop gang + for (j = 0; j < 32; j++) + { + res = 0; + + #pragma acc loop worker reduction(+:res) + for (i = 0; i < 32; i++) + res += arr[j * 32 + i]; + + out[j] = res; + } + } + + for (j = 0; j < 32; j++) + { + hres = 0; + + for (i = 0; i < 32; i++) + hres += arr[j * 32 + i]; + + assert (out[j] == hres); + } +} + + +/* Test of reduction on loop directive (workers and vectors, private reduction + variable). */ + +void wv_p_1() +{ + int i, j, arr[1024], out[32], res = 0, hres = 0; + + for (i = 0; i < 1024; i++) + arr[i] = i; + + #pragma acc parallel num_gangs(32) num_workers(32) vector_length(32) \ + private(res) copyout(out) + { + #pragma acc loop gang + for (j = 0; j < 32; j++) + { + res = 0; + + #pragma acc loop worker vector reduction(+:res) + for (i = 0; i < 32; i++) + res += arr[j * 32 + i]; + + out[j] = res; + } + } + + for (j = 0; j < 32; j++) + { + hres = 0; + + for (i = 0; i < 32; i++) + hres += arr[j * 32 + i]; + + assert (out[j] == hres); + } +} + + +/* Test of reduction on loop directive (workers and vectors, private reduction + variable). */ + +void wv_p_2() +{ + int i, j, arr[32768], out[32], 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) \ + private(res) copyout(out) + { + #pragma acc loop gang + for (j = 0; j < 32; j++) + { + res = j; + + #pragma acc loop worker reduction(+:res) + for (i = 0; i < 1024; i++) + res += arr[j * 1024 + i]; + + #pragma acc loop vector reduction(+:res) + for (i = 1023; i >= 0; i--) + res += arr[j * 1024 + i]; + + out[j] = res; + } + } + + for (j = 0; j < 32; j++) + { + hres = j; + + for (i = 0; i < 1024; i++) + hres += arr[j * 1024 + i] * 2; + + assert (out[j] == hres); + } +} + + +/* Test of reduction on loop directive (workers and vectors, private reduction + variable: gang-redundant mode). */ + +void wv_p_3() +{ + int i, arr[1024], out[32], res = 0, hres = 0; + + for (i = 0; i < 1024; i++) + arr[i] = i ^ 33; + + #pragma acc parallel num_gangs(32) num_workers(32) vector_length(32) \ + private(res) copyin(arr) copyout(out) + { + /* Private variables aren't initialized by default in openacc. */ + res = 0; + + /* "res" should be available at the end of the following loop (and should + have the same value redundantly in each gang). */ + #pragma acc loop worker vector reduction(+:res) + for (i = 0; i < 1024; i++) + res += arr[i]; + + #pragma acc loop gang (static: 1) + for (i = 0; i < 32; i++) + out[i] = res; + } + + for (i = 0; i < 1024; i++) + hres += arr[i]; + + for (i = 0; i < 32; i++) + assert (out[i] == hres); +} + + int main() { g_np_1(); gv_np_1(); gw_np_1(); gwv_np_1(); + gwv_np_2(); + gwv_np_3(); + gwv_np_4(); + v_p_1(); + v_p_2(); + w_p_1(); + wv_p_1(); + wv_p_2(); + wv_p_3(); return 0; }