From patchwork Wed Nov 4 20:51:55 2015 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Nathan Sidwell X-Patchwork-Id: 540174 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 4E3E7140B0F for ; Thu, 5 Nov 2015 07:52:11 +1100 (AEDT) Authentication-Results: ozlabs.org; dkim=pass (1024-bit key; unprotected) header.d=gcc.gnu.org header.i=@gcc.gnu.org header.b=FuZb30F4; 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:to :from:subject:message-id:date:mime-version:content-type; q=dns; s=default; b=pcJdKuQl4v4zgoR1KQd8/LI0cnTK7+tZ+k89+m46p9Pt1b9rgF AGsFZX0lCuDeScnWhh8zCpdBum94jjMAnKVFyWncBfKyNqYKONRFYlBMUc3qsfsE EWCsKVO6zBI3smAXK4DBC5LH9beTHyTXmMsQEXcHCluXDXQG+YHts5hU8= 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:to :from:subject:message-id:date:mime-version:content-type; s= default; bh=ai1bS2sUpv05oyPH9/AgSe4nfaA=; b=FuZb30F4C8BK9GHUWA4U CZH6vcpPgXVNKTFFPaVsCA1a2nEnfphYMByHVs7IyIEtwb6rRTIe4e8OVd0hTgXk qFPWyF0CeBUe28MpuTiYJ0eRPbR2PfvORgIeZPY8qYtXSkk4yTsocHKiF7y2pd+q 6O2jCXm5FUpoYXQoyJ93pAc= Received: (qmail 27116 invoked by alias); 4 Nov 2015 20:52:02 -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 26587 invoked by uid 89); 4 Nov 2015 20:52:02 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=0.9 required=5.0 tests=BAYES_50, FREEMAIL_FROM, KAM_ASCII_DIVIDERS, RCVD_IN_DNSWL_LOW, SPF_PASS autolearn=no version=3.3.2 X-HELO: mail-yk0-f174.google.com Received: from mail-yk0-f174.google.com (HELO mail-yk0-f174.google.com) (209.85.160.174) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with (AES128-GCM-SHA256 encrypted) ESMTPS; Wed, 04 Nov 2015 20:51:59 +0000 Received: by ykdr3 with SMTP id r3so94930527ykd.1 for ; Wed, 04 Nov 2015 12:51:57 -0800 (PST) X-Received: by 10.31.8.201 with SMTP id 192mr3786612vki.67.1446670317067; Wed, 04 Nov 2015 12:51:57 -0800 (PST) Received: from ?IPv6:2601:181:c000:c497:a2a8:cdff:fe3e:b48? ([2601:181:c000:c497:a2a8:cdff:fe3e:b48]) by smtp.googlemail.com with ESMTPSA id e192sm2105301vkd.18.2015.11.04.12.51.56 (version=TLSv1.2 cipher=ECDHE-RSA-AES128-GCM-SHA256 bits=128/128); Wed, 04 Nov 2015 12:51:56 -0800 (PST) To: GCC Patches From: Nathan Sidwell Subject: OpenaCC dimension checking Message-ID: <563A6FEB.8070300@acm.org> Date: Wed, 4 Nov 2015 15:51:55 -0500 User-Agent: Mozilla/5.0 (X11; Linux x86_64; rv:38.0) Gecko/20100101 Thunderbird/38.3.0 MIME-Version: 1.0 Now the core of the execution model is committed, I've applied this patch to enable the nvptx dimension checking. We force the vector size to be 32 in all cases, and check the worker size is not above 32. Warnings are given if the user specifies an unacceptable dimension. This patch exposed some unacceptable testcases which I've modified to specify an acceptable vector length. I also took the opportunity to fix up their reduction variable copying, which is only working right now because we default to copy, rather than firstprvate (that's the next patch for review). Also a new testcase to exercise the error handling. Committed to trunk. nathan 2015-11-04 Nathan Sidwell gcc/ * config/nvptx/nvptx.c (nvptx_goacc_validate_dims): Add checking. libgomp/ * testsuite/libgomp.oacc-fortran/reduction-1.f90: Fix dimensions and reduction copy. * testsuite/libgomp.oacc-fortran/reduction-2.f90: Likewise. * testsuite/libgomp.oacc-fortran/reduction-3.f90: Likewise. * testsuite/libgomp.oacc-fortran/reduction-4.f90: Likewise. * testsuite/libgomp.oacc-fortran/reduction-6.f90: Likewise. * testsuite/libgomp.oacc-c-c++-common/par-reduction-1.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/reduction-3.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/collapse-2.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/par-reduction-2.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/reduction-4.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/reduction-initial-1.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/reduction-1.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/reduction-5.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/reduction-2.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/parallel-dims.c: New. Index: gcc/config/nvptx/nvptx.c =================================================================== --- gcc/config/nvptx/nvptx.c (revision 229771) +++ gcc/config/nvptx/nvptx.c (working copy) @@ -3472,8 +3472,29 @@ nvptx_goacc_validate_dims (tree ARG_UNUS { bool changed = false; - /* TODO: Leave dimensions unaltered. Reductions need - porting before filtering dimensions makes sense. */ + /* The vector size must be 32, unless this is a SEQ routine. */ + if (fn_level <= GOMP_DIM_VECTOR + && dims[GOMP_DIM_VECTOR] != PTX_VECTOR_LENGTH) + { + if (dims[GOMP_DIM_VECTOR] >= 0 && fn_level < 0) + warning_at (DECL_SOURCE_LOCATION (decl), 0, + dims[GOMP_DIM_VECTOR] + ? "using vector_length (%d), ignoring %d" + : "using vector_length (%d), ignoring runtime setting", + PTX_VECTOR_LENGTH, dims[GOMP_DIM_VECTOR]); + dims[GOMP_DIM_VECTOR] = PTX_VECTOR_LENGTH; + changed = true; + } + + /* Check the num workers is not too large. */ + if (dims[GOMP_DIM_WORKER] > PTX_WORKER_LENGTH) + { + warning_at (DECL_SOURCE_LOCATION (decl), 0, + "using num_workers (%d), ignoring %d", + PTX_WORKER_LENGTH, dims[GOMP_DIM_WORKER]); + dims[GOMP_DIM_WORKER] = PTX_WORKER_LENGTH; + changed = true; + } return changed; } Index: libgomp/testsuite/libgomp.oacc-fortran/reduction-1.f90 =================================================================== --- libgomp/testsuite/libgomp.oacc-fortran/reduction-1.f90 (revision 229771) +++ libgomp/testsuite/libgomp.oacc-fortran/reduction-1.f90 (working copy) @@ -5,7 +5,7 @@ program reduction_1 implicit none - integer, parameter :: n = 10, vl = 2 + integer, parameter :: n = 10, vl = 32 integer :: i, vresult, result logical :: lresult, lvresult integer, dimension (n) :: array @@ -19,7 +19,7 @@ program reduction_1 ! '+' reductions - !$acc parallel vector_length(vl) num_gangs(1) + !$acc parallel vector_length(vl) num_gangs(1) copy(result) !$acc loop reduction(+:result) do i = 1, n result = result + array(i) @@ -38,7 +38,7 @@ program reduction_1 ! '*' reductions - !$acc parallel vector_length(vl) num_gangs(1) + !$acc parallel vector_length(vl) num_gangs(1) copy(result) !$acc loop reduction(*:result) do i = 1, n result = result * array(i) @@ -57,7 +57,7 @@ program reduction_1 ! 'max' reductions - !$acc parallel vector_length(vl) num_gangs(1) + !$acc parallel vector_length(vl) num_gangs(1) copy(result) !$acc loop reduction(max:result) do i = 1, n result = max (result, array(i)) @@ -76,7 +76,7 @@ program reduction_1 ! 'min' reductions - !$acc parallel vector_length(vl) num_gangs(1) + !$acc parallel vector_length(vl) num_gangs(1) copy(result) !$acc loop reduction(min:result) do i = 1, n result = min (result, array(i)) @@ -95,7 +95,7 @@ program reduction_1 ! 'iand' reductions - !$acc parallel vector_length(vl) num_gangs(1) + !$acc parallel vector_length(vl) num_gangs(1) copy(result) !$acc loop reduction(iand:result) do i = 1, n result = iand (result, array(i)) @@ -114,7 +114,7 @@ program reduction_1 ! 'ior' reductions - !$acc parallel vector_length(vl) num_gangs(1) + !$acc parallel vector_length(vl) num_gangs(1) copy(result) !$acc loop reduction(ior:result) do i = 1, n result = ior (result, array(i)) @@ -133,7 +133,7 @@ program reduction_1 ! 'ieor' reductions - !$acc parallel vector_length(vl) num_gangs(1) + !$acc parallel vector_length(vl) num_gangs(1) copy(result) !$acc loop reduction(ieor:result) do i = 1, n result = ieor (result, array(i)) @@ -152,7 +152,7 @@ program reduction_1 ! '.and.' reductions - !$acc parallel vector_length(vl) num_gangs(1) + !$acc parallel vector_length(vl) num_gangs(1) copy(lresult) !$acc loop reduction(.and.:lresult) do i = 1, n lresult = lresult .and. (array(i) .ge. 5) @@ -171,7 +171,7 @@ program reduction_1 ! '.or.' reductions - !$acc parallel vector_length(vl) num_gangs(1) + !$acc parallel vector_length(vl) num_gangs(1) copy(lresult) !$acc loop reduction(.or.:lresult) do i = 1, n lresult = lresult .or. (array(i) .ge. 5) @@ -190,7 +190,7 @@ program reduction_1 ! '.eqv.' reductions - !$acc parallel vector_length(vl) num_gangs(1) + !$acc parallel vector_length(vl) num_gangs(1) copy(lresult) !$acc loop reduction(.eqv.:lresult) do i = 1, n lresult = lresult .eqv. (array(i) .ge. 5) @@ -209,7 +209,7 @@ program reduction_1 ! '.neqv.' reductions - !$acc parallel vector_length(vl) num_gangs(1) + !$acc parallel vector_length(vl) num_gangs(1) copy(lresult) !$acc loop reduction(.neqv.:lresult) do i = 1, n lresult = lresult .neqv. (array(i) .ge. 5) Index: libgomp/testsuite/libgomp.oacc-fortran/reduction-2.f90 =================================================================== --- libgomp/testsuite/libgomp.oacc-fortran/reduction-2.f90 (revision 229771) +++ libgomp/testsuite/libgomp.oacc-fortran/reduction-2.f90 (working copy) @@ -5,7 +5,7 @@ program reduction_2 implicit none - integer, parameter :: n = 10, vl = 2 + integer, parameter :: n = 10, vl = 32 integer :: i real, parameter :: e = .001 real :: vresult, result @@ -21,7 +21,7 @@ program reduction_2 ! '+' reductions - !$acc parallel vector_length(vl) num_gangs(1) + !$acc parallel vector_length(vl) num_gangs(1) copy(result) !$acc loop reduction(+:result) do i = 1, n result = result + array(i) @@ -40,7 +40,7 @@ program reduction_2 ! '*' reductions - !$acc parallel vector_length(vl) num_gangs(1) + !$acc parallel vector_length(vl) num_gangs(1) copy(result) !$acc loop reduction(*:result) do i = 1, n result = result * array(i) @@ -59,7 +59,7 @@ program reduction_2 ! 'max' reductions - !$acc parallel vector_length(vl) num_gangs(1) + !$acc parallel vector_length(vl) num_gangs(1) copy(result) !$acc loop reduction(max:result) do i = 1, n result = max (result, array(i)) @@ -78,7 +78,7 @@ program reduction_2 ! 'min' reductions - !$acc parallel vector_length(vl) num_gangs(1) + !$acc parallel vector_length(vl) num_gangs(1) copy(result) !$acc loop reduction(min:result) do i = 1, n result = min (result, array(i)) @@ -97,7 +97,7 @@ program reduction_2 ! '.and.' reductions - !$acc parallel vector_length(vl) num_gangs(1) + !$acc parallel vector_length(vl) num_gangs(1) copy(lresult) !$acc loop reduction(.and.:lresult) do i = 1, n lresult = lresult .and. (array(i) .ge. 5) @@ -116,7 +116,7 @@ program reduction_2 ! '.or.' reductions - !$acc parallel vector_length(vl) num_gangs(1) + !$acc parallel vector_length(vl) num_gangs(1) copy(lresult) !$acc loop reduction(.or.:lresult) do i = 1, n lresult = lresult .or. (array(i) .ge. 5) @@ -135,7 +135,7 @@ program reduction_2 ! '.eqv.' reductions - !$acc parallel vector_length(vl) num_gangs(1) + !$acc parallel vector_length(vl) num_gangs(1) copy(lresult) !$acc loop reduction(.eqv.:lresult) do i = 1, n lresult = lresult .eqv. (array(i) .ge. 5) @@ -154,7 +154,7 @@ program reduction_2 ! '.neqv.' reductions - !$acc parallel vector_length(vl) num_gangs(1) + !$acc parallel vector_length(vl) num_gangs(1) copy(lresult) !$acc loop reduction(.neqv.:lresult) do i = 1, n lresult = lresult .neqv. (array(i) .ge. 5) Index: libgomp/testsuite/libgomp.oacc-fortran/reduction-3.f90 =================================================================== --- libgomp/testsuite/libgomp.oacc-fortran/reduction-3.f90 (revision 229771) +++ libgomp/testsuite/libgomp.oacc-fortran/reduction-3.f90 (working copy) @@ -5,7 +5,7 @@ program reduction_3 implicit none - integer, parameter :: n = 10, vl = 2 + integer, parameter :: n = 10, vl = 32 integer :: i double precision, parameter :: e = .001 double precision :: vresult, result @@ -21,7 +21,7 @@ program reduction_3 ! '+' reductions - !$acc parallel vector_length(vl) num_gangs(1) + !$acc parallel vector_length(vl) num_gangs(1) copy(result) !$acc loop reduction(+:result) do i = 1, n result = result + array(i) @@ -40,7 +40,7 @@ program reduction_3 ! '*' reductions - !$acc parallel vector_length(vl) num_gangs(1) + !$acc parallel vector_length(vl) num_gangs(1) copy(result) !$acc loop reduction(*:result) do i = 1, n result = result * array(i) @@ -59,7 +59,7 @@ program reduction_3 ! 'max' reductions - !$acc parallel vector_length(vl) num_gangs(1) + !$acc parallel vector_length(vl) num_gangs(1) copy(result) !$acc loop reduction(max:result) do i = 1, n result = max (result, array(i)) @@ -78,7 +78,7 @@ program reduction_3 ! 'min' reductions - !$acc parallel vector_length(vl) num_gangs(1) + !$acc parallel vector_length(vl) num_gangs(1) copy(result) !$acc loop reduction(min:result) do i = 1, n result = min (result, array(i)) @@ -97,7 +97,7 @@ program reduction_3 ! '.and.' reductions - !$acc parallel vector_length(vl) num_gangs(1) + !$acc parallel vector_length(vl) num_gangs(1) copy(lresult) !$acc loop reduction(.and.:lresult) do i = 1, n lresult = lresult .and. (array(i) .ge. 5) @@ -116,7 +116,7 @@ program reduction_3 ! '.or.' reductions - !$acc parallel vector_length(vl) num_gangs(1) + !$acc parallel vector_length(vl) num_gangs(1) copy(lresult) !$acc loop reduction(.or.:lresult) do i = 1, n lresult = lresult .or. (array(i) .ge. 5) @@ -135,7 +135,7 @@ program reduction_3 ! '.eqv.' reductions - !$acc parallel vector_length(vl) num_gangs(1) + !$acc parallel vector_length(vl) num_gangs(1) copy(lresult) !$acc loop reduction(.eqv.:lresult) do i = 1, n lresult = lresult .eqv. (array(i) .ge. 5) @@ -154,7 +154,7 @@ program reduction_3 ! '.neqv.' reductions - !$acc parallel vector_length(vl) num_gangs(1) + !$acc parallel vector_length(vl) num_gangs(1) copy(lresult) !$acc loop reduction(.neqv.:lresult) do i = 1, n lresult = lresult .neqv. (array(i) .ge. 5) Index: libgomp/testsuite/libgomp.oacc-fortran/reduction-4.f90 =================================================================== --- libgomp/testsuite/libgomp.oacc-fortran/reduction-4.f90 (revision 229771) +++ libgomp/testsuite/libgomp.oacc-fortran/reduction-4.f90 (working copy) @@ -19,7 +19,7 @@ program reduction_4 ! '+' reductions - !$acc parallel vector_length(vl) num_gangs(1) + !$acc parallel vector_length(vl) num_gangs(1) copy(result) !$acc loop reduction(+:result) do i = 1, n result = result + array(i) Index: libgomp/testsuite/libgomp.oacc-fortran/reduction-6.f90 =================================================================== --- libgomp/testsuite/libgomp.oacc-fortran/reduction-6.f90 (revision 229771) +++ libgomp/testsuite/libgomp.oacc-fortran/reduction-6.f90 (working copy) @@ -11,7 +11,7 @@ program reduction vs1 = 0 vs2 = 0 - !$acc parallel vector_length (32) + !$acc parallel vector_length (32) copy(s1, s2) !$acc loop reduction(+:s1, s2) do i = 1, n s1 = s1 + 1 Index: libgomp/testsuite/libgomp.oacc-c-c++-common/par-reduction-1.c =================================================================== --- libgomp/testsuite/libgomp.oacc-c-c++-common/par-reduction-1.c (revision 229771) +++ libgomp/testsuite/libgomp.oacc-c-c++-common/par-reduction-1.c (working copy) @@ -10,8 +10,7 @@ main (int argc, char *argv[]) #else # define GANGS 256 #endif - #pragma acc parallel num_gangs(GANGS) num_workers(1) vector_length(1) \ - copy(res2) + #pragma acc parallel num_gangs(GANGS) copy(res2) { #pragma acc atomic res2 += 5; @@ -28,8 +27,7 @@ main (int argc, char *argv[]) #else # define GANGS 8 #endif - #pragma acc parallel num_gangs(GANGS) num_workers(1) vector_length(1) \ - copy(res2) + #pragma acc parallel num_gangs(GANGS) copy(res2) { #pragma acc atomic res2 *= 5; Index: libgomp/testsuite/libgomp.oacc-c-c++-common/reduction-3.c =================================================================== --- libgomp/testsuite/libgomp.oacc-c-c++-common/reduction-3.c (revision 229771) +++ libgomp/testsuite/libgomp.oacc-c-c++-common/reduction-3.c (working copy) @@ -23,7 +23,7 @@ main(void) vresult = 0; /* '+' reductions. */ -#pragma acc parallel vector_length (vl) +#pragma acc parallel vector_length (vl) copy(result) #pragma acc loop reduction (+:result) for (i = 0; i < n; i++) result += array[i]; @@ -39,7 +39,7 @@ main(void) vresult = 0; /* '*' reductions. */ -#pragma acc parallel vector_length (vl) +#pragma acc parallel vector_length (vl) copy(result) #pragma acc loop reduction (*:result) for (i = 0; i < n; i++) result *= array[i]; @@ -91,7 +91,7 @@ main(void) lvresult = false; /* '&&' reductions. */ -#pragma acc parallel vector_length (vl) +#pragma acc parallel vector_length (vl) copy(lresult) #pragma acc loop reduction (&&:lresult) for (i = 0; i < n; i++) lresult = lresult && (result > array[i]); @@ -110,7 +110,7 @@ main(void) lvresult = false; /* '||' reductions. */ -#pragma acc parallel vector_length (vl) +#pragma acc parallel vector_length (vl) copy(lresult) #pragma acc loop reduction (||:lresult) for (i = 0; i < n; i++) lresult = lresult || (result > array[i]); Index: libgomp/testsuite/libgomp.oacc-c-c++-common/collapse-2.c =================================================================== --- libgomp/testsuite/libgomp.oacc-c-c++-common/collapse-2.c (revision 229771) +++ libgomp/testsuite/libgomp.oacc-c-c++-common/collapse-2.c (working copy) @@ -8,7 +8,7 @@ main (void) int i, j, k, l = 0, f = 0, x = 0; int m1 = 4, m2 = -5, m3 = 17; - #pragma acc parallel +#pragma acc parallel copy(l) #pragma acc loop collapse(3) reduction(+:l) for (i = -2; i < m1; i++) for (j = m2; j < -2; j++) Index: libgomp/testsuite/libgomp.oacc-c-c++-common/par-reduction-2.c =================================================================== --- libgomp/testsuite/libgomp.oacc-c-c++-common/par-reduction-2.c (revision 229771) +++ libgomp/testsuite/libgomp.oacc-c-c++-common/par-reduction-2.c (working copy) @@ -11,8 +11,7 @@ main (int argc, char *argv[]) #else # define GANGS 256 #endif - #pragma acc parallel num_gangs(GANGS) num_workers(1) vector_length(1) \ - copy(res2) async(1) + #pragma acc parallel num_gangs(GANGS) copy(res2) async(1) { #pragma acc atomic res2 += 5; @@ -31,8 +30,7 @@ main (int argc, char *argv[]) #else # define GANGS 8 #endif - #pragma acc parallel num_gangs(GANGS) num_workers(1) vector_length(1) \ - copy(res2) async(1) + #pragma acc parallel num_gangs(GANGS) copy(res2) async(1) { #pragma acc atomic res2 *= 5; Index: libgomp/testsuite/libgomp.oacc-c-c++-common/reduction-4.c =================================================================== --- libgomp/testsuite/libgomp.oacc-c-c++-common/reduction-4.c (revision 229771) +++ libgomp/testsuite/libgomp.oacc-c-c++-common/reduction-4.c (working copy) @@ -24,7 +24,7 @@ main(void) vresult = 0; /* '+' reductions. */ -#pragma acc parallel vector_length (vl) +#pragma acc parallel vector_length (vl) copy(result) #pragma acc loop reduction (+:result) for (i = 0; i < n; i++) result += array[i]; @@ -94,7 +94,7 @@ main(void) lvresult = false; /* '&&' reductions. */ -#pragma acc parallel vector_length (vl) +#pragma acc parallel vector_length (vl) copy(lresult) #pragma acc loop reduction (&&:lresult) for (i = 0; i < n; i++) lresult = lresult && (creal(result) > creal(array[i])); @@ -113,7 +113,7 @@ main(void) lvresult = false; /* '||' reductions. */ -#pragma acc parallel vector_length (vl) +#pragma acc parallel vector_length (vl) copy(lresult) #pragma acc loop reduction (||:lresult) for (i = 0; i < n; i++) lresult = lresult || (creal(result) > creal(array[i])); Index: libgomp/testsuite/libgomp.oacc-c-c++-common/reduction-initial-1.c =================================================================== --- libgomp/testsuite/libgomp.oacc-c-c++-common/reduction-initial-1.c (revision 229771) +++ libgomp/testsuite/libgomp.oacc-c-c++-common/reduction-initial-1.c (working copy) @@ -4,13 +4,13 @@ int main(void) { #define I 5 -#define N 11 +#define N 32 #define A 8 int a = A; int s = I; -#pragma acc parallel vector_length(N) +#pragma acc parallel vector_length(N) copy(s) { int i; #pragma acc loop reduction(+:s) Index: libgomp/testsuite/libgomp.oacc-c-c++-common/reduction-1.c =================================================================== --- libgomp/testsuite/libgomp.oacc-c-c++-common/reduction-1.c (revision 229771) +++ libgomp/testsuite/libgomp.oacc-c-c++-common/reduction-1.c (working copy) @@ -13,7 +13,7 @@ { \ type res, vres; \ res = (init); \ -DO_PRAGMA (acc parallel vector_length (vl))\ + DO_PRAGMA (acc parallel vector_length (vl) copy(res)) \ DO_PRAGMA (acc loop reduction (op:res))\ for (i = 0; i < n; i++) \ res = res op (b); \ @@ -63,7 +63,7 @@ test_reductions_bool (void) { \ type res, vres; \ res = (init); \ -DO_PRAGMA (acc parallel vector_length (vl))\ +DO_PRAGMA (acc parallel vector_length (vl) copy(res))\ DO_PRAGMA (acc loop reduction (op:res))\ for (i = 0; i < n; i++) \ res = op (res, (b)); \ Index: libgomp/testsuite/libgomp.oacc-c-c++-common/reduction-5.c =================================================================== --- libgomp/testsuite/libgomp.oacc-c-c++-common/reduction-5.c (revision 229771) +++ libgomp/testsuite/libgomp.oacc-c-c++-common/reduction-5.c (working copy) @@ -8,7 +8,7 @@ main (void) int n = 100; int i; -#pragma acc parallel vector_length (32) +#pragma acc parallel vector_length (32) copy(s1,s2) #pragma acc loop reduction (+:s1, s2) for (i = 0; i < n; i++) { Index: libgomp/testsuite/libgomp.oacc-c-c++-common/reduction-2.c =================================================================== --- libgomp/testsuite/libgomp.oacc-c-c++-common/reduction-2.c (revision 229771) +++ libgomp/testsuite/libgomp.oacc-c-c++-common/reduction-2.c (working copy) @@ -23,7 +23,7 @@ main(void) vresult = 0; /* '+' reductions. */ -#pragma acc parallel vector_length (vl) +#pragma acc parallel vector_length (vl) copy(result) #pragma acc loop reduction (+:result) for (i = 0; i < n; i++) result += array[i]; @@ -39,7 +39,7 @@ main(void) vresult = 0; /* '*' reductions. */ -#pragma acc parallel vector_length (vl) +#pragma acc parallel vector_length (vl) copy(result) #pragma acc loop reduction (*:result) for (i = 0; i < n; i++) result *= array[i]; @@ -91,7 +91,7 @@ main(void) lvresult = false; /* '&&' reductions. */ -#pragma acc parallel vector_length (vl) +#pragma acc parallel vector_length (vl) copy(lresult) #pragma acc loop reduction (&&:lresult) for (i = 0; i < n; i++) lresult = lresult && (result > array[i]); @@ -110,7 +110,7 @@ main(void) lvresult = false; /* '||' reductions. */ -#pragma acc parallel vector_length (vl) +#pragma acc parallel vector_length (vl) copy(lresult) #pragma acc loop reduction (||:lresult) for (i = 0; i < n; i++) lresult = lresult || (result > array[i]); Index: libgomp/testsuite/libgomp.oacc-c-c++-common/parallel-dims.c =================================================================== --- libgomp/testsuite/libgomp.oacc-c-c++-common/parallel-dims.c (revision 0) +++ libgomp/testsuite/libgomp.oacc-c-c++-common/parallel-dims.c (working copy) @@ -0,0 +1,17 @@ +/* { dg-do run { target openacc_nvidia_accel_selected } } */ + +/* Worker and vector size checks. Picked an outrageously large + value. */ + +int main () +{ +#pragma acc parallel num_workers (2<<20) /* { dg-error "using num_workers" } */ + { + } + +#pragma acc parallel vector_length (2<<20) /* { dg-error "using vector_length" } */ + { + } + + return 0; +}