From patchwork Fri Jan 21 17:48:20 2022 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 8bit X-Patchwork-Submitter: Thomas Schwinge X-Patchwork-Id: 50329 Return-Path: X-Original-To: patchwork@sourceware.org Delivered-To: patchwork@sourceware.org Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id 695A1385AC19 for ; Fri, 21 Jan 2022 17:48:49 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from esa3.mentor.iphmx.com (esa3.mentor.iphmx.com [68.232.137.180]) by sourceware.org (Postfix) with ESMTPS id 56CCE3858423 for ; Fri, 21 Jan 2022 17:48:28 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.1 sourceware.org 56CCE3858423 Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=mentor.com IronPort-SDR: EiPHRHw+2mNfHcsNXZZm7jnsjhhkstGBA6Mc/3bTRVlad/uBzeYrFAw2iKqr01a4y8ERvFCZwN 23tcjCAYHX0UP3jbEDH6SDmCbUycEfFaQ/DrRf93bLIK/ec7JzJwVEX698O8L9nnXkYb4IjcIn biTXtpg9wT2/s4iKQNupwESlgZxlsdFgWSTeIn7Gop1vH2ZRkGw1A3O7TD3ScVU3FdZDXkTc7s k7aIuR9ilonWan4+RajWXTIQcZF19SJ6KhXGxC71l7WLeoE9b8ZOATA85FZWNzfBz2yDkdp1Sh orJ3q6wNasQxx/oUMAQc92fd X-IronPort-AV: E=Sophos;i="5.88,306,1635235200"; d="scan'208,223";a="70925703" Received: from orw-gwy-02-in.mentorg.com ([192.94.38.167]) by esa3.mentor.iphmx.com with ESMTP; 21 Jan 2022 09:48:27 -0800 IronPort-SDR: c9BKa212KnxJvXZlptFpvSgqAz+ffx37GMs3zAksNjrqZi+nqew4jF9NJDfsVFGZZB0dsC/kJA xOVTzbKtmi2/71LWJKmNJ4aDmwKPGvNlwICn/6VPiZPkqmd6AAvh8MP8SbjbvaFSbVaGihxNia UEtSrVPXKdPzDPCNAFbtZlsKv2yJZaF5DNs7ekTe3gwKj8tqsz1CO+KZYXgrKGSS8rdhvUdfEV ERloBTmBMmpnDAFvyYkrXJ9Y+E+AV5YsVM4klI+ec7EmbEhwV7RTJKY7AoagLRQSX9cHbhfPbF MiE= From: Thomas Schwinge To: Subject: Strengthen a few OpenACC test cases User-Agent: Notmuch/0.29.3+94~g74c3f1b (https://notmuchmail.org) Emacs/27.1 (x86_64-pc-linux-gnu) Date: Fri, 21 Jan 2022 18:48:20 +0100 Message-ID: <871r11c74r.fsf@euler.schwinge.homeip.net> MIME-Version: 1.0 X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: svr-ies-mbx-05.mgc.mentorg.com (139.181.222.5) To svr-ies-mbx-01.mgc.mentorg.com (139.181.222.1) X-Spam-Status: No, score=-12.1 required=5.0 tests=BAYES_00, GIT_PATCH_0, HEADER_FROM_DIFFERENT_DOMAINS, KAM_DMARC_STATUS, SPF_HELO_PASS, SPF_PASS, TXREP autolearn=ham autolearn_force=no version=3.4.4 X-Spam-Checker-Version: SpamAssassin 3.4.4 (2020-01-24) on server2.sourceware.org X-BeenThere: gcc-patches@gcc.gnu.org X-Mailman-Version: 2.1.29 Precedence: list List-Id: Gcc-patches mailing list List-Unsubscribe: , List-Archive: List-Post: List-Help: List-Subscribe: , Errors-To: gcc-patches-bounces+patchwork=sourceware.org@gcc.gnu.org Sender: "Gcc-patches" Hi! Pushed to master branch commit 087e545747ca9ee977e84326877b0ce1bc4c383a "Strengthen a few OpenACC test cases", see attached. Grüße Thomas ----------------- Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955 From 087e545747ca9ee977e84326877b0ce1bc4c383a Mon Sep 17 00:00:00 2001 From: Thomas Schwinge Date: Fri, 21 Jan 2022 12:48:28 +0100 Subject: [PATCH] Strengthen a few OpenACC test cases Rather than rubber-stamp whatever requested vs. actual device kernel launch configuration happens, actually (again) verify the requested values (modulo expected variations). This better highlights that "AMD GCN has an upper limit of 'num_workers(16)'", and the deficiency that "AMD GCN uses the autovectorizer for the vector dimension: the use of a function call in vector-partitioned code [...] is not currently supported". And, this removes several instances of race conditions, where variables are concurrently written to in OpenACC gang-redundant mode. libgomp/ * testsuite/libgomp.oacc-c-c++-common/loop-gwv-1.c: Strengthen. * testsuite/libgomp.oacc-c-c++-common/loop-gwv-2.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/loop-red-gwv-1.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/loop-red-v-1.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/loop-red-v-2.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/loop-red-w-1.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/loop-red-w-2.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/loop-red-wv-1.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/loop-v-1.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/loop-w-1.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/loop-wv-1.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/routine-gwv-1.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/routine-v-1.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/routine-w-1.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/routine-wv-1.c: Likewise. --- .../libgomp.oacc-c-c++-common/loop-gwv-1.c | 25 ++++++++++---- .../libgomp.oacc-c-c++-common/loop-gwv-2.c | 34 +++++++++++++++---- .../loop-red-gwv-1.c | 22 +++++++++--- .../libgomp.oacc-c-c++-common/loop-red-v-1.c | 11 ++++-- .../libgomp.oacc-c-c++-common/loop-red-v-2.c | 11 ++++-- .../libgomp.oacc-c-c++-common/loop-red-w-1.c | 13 +++++-- .../libgomp.oacc-c-c++-common/loop-red-w-2.c | 13 +++++-- .../libgomp.oacc-c-c++-common/loop-red-wv-1.c | 19 ++++++++--- .../libgomp.oacc-c-c++-common/loop-v-1.c | 13 +++++-- .../libgomp.oacc-c-c++-common/loop-w-1.c | 13 +++++-- .../libgomp.oacc-c-c++-common/loop-wv-1.c | 19 ++++++++--- .../libgomp.oacc-c-c++-common/routine-gwv-1.c | 21 +++++++++--- .../libgomp.oacc-c-c++-common/routine-v-1.c | 13 +++++-- .../libgomp.oacc-c-c++-common/routine-w-1.c | 13 +++++-- .../libgomp.oacc-c-c++-common/routine-wv-1.c | 19 ++++++++--- 15 files changed, 202 insertions(+), 57 deletions(-) diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-gwv-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-gwv-1.c index e5ed2ab7006..d3f6ea24e7e 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-gwv-1.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-gwv-1.c @@ -19,9 +19,12 @@ int main () for (ix = 0; ix < N;ix++) ary[ix] = -1; - -#pragma acc parallel num_gangs(32) num_workers(32) vector_length(32) \ - copy(ary) copy(ondev) copyout(gangsize, workersize, vectorsize) + +#define NG 32 +#define NW 32 +#define VL 32 +#pragma acc parallel num_gangs(NG) num_workers(NW) vector_length(VL) \ + copy(ary) copy(ondev) /* { dg-note {variable 'ix' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-2 } */ { #pragma acc loop gang worker vector @@ -45,11 +48,19 @@ int main () else ary[ix] = ix; } - - gangsize = __builtin_goacc_parlevel_size (GOMP_DIM_GANG); - workersize = __builtin_goacc_parlevel_size (GOMP_DIM_WORKER); - vectorsize = __builtin_goacc_parlevel_size (GOMP_DIM_VECTOR); } + gangsize = NG; + workersize = NW; + vectorsize = VL; +#ifdef ACC_DEVICE_TYPE_radeon + /* AMD GCN has an upper limit of 'num_workers(16)'. */ + if (workersize > 16) + workersize = 16; + /* AMD GCN uses the autovectorizer for the vector dimension: the use + of a function call in vector-partitioned code in this test is not + currently supported. */ + vectorsize = 1; +#endif for (ix = 0; ix < N; ix++) { diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-gwv-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-gwv-2.c index e73ed6064eb..4b761f0f624 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-gwv-2.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-gwv-2.c @@ -46,14 +46,17 @@ int main () int ary[N]; int ix; int exit = 0; - int gangsize = 0, workersize = 0, vectorsize = 0; + int gangsize, workersize, vectorsize; int *gangdist, *workerdist, *vectordist; for (ix = 0; ix < N;ix++) ary[ix] = -1; -#pragma acc parallel num_gangs(32) num_workers(32) vector_length(32) \ - copy(ary) copyout(gangsize, workersize, vectorsize) +#define NG 32 +#define NW 32 +#define VL 32 +#pragma acc parallel num_gangs(NG) num_workers(NW) vector_length(VL) \ + copy(ary) /* { dg-note {variable 'ix' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-2 } */ { #pragma acc loop gang worker vector @@ -71,11 +74,23 @@ int main () ary[ix] = (g << 16) | (w << 8) | v; } - - gangsize = __builtin_goacc_parlevel_size (GOMP_DIM_GANG); - workersize = __builtin_goacc_parlevel_size (GOMP_DIM_WORKER); - vectorsize = __builtin_goacc_parlevel_size (GOMP_DIM_VECTOR); } + gangsize = NG; + workersize = NW; + vectorsize = VL; +#if defined ACC_DEVICE_TYPE_host + gangsize = 1; + workersize = 1; + vectorsize = 1; +#elif defined ACC_DEVICE_TYPE_radeon + /* AMD GCN has an upper limit of 'num_workers(16)'. */ + if (workersize > 16) + workersize = 16; + /* AMD GCN uses the autovectorizer for the vector dimension: the use + of a function call in vector-partitioned code in this test is not + currently supported. */ + vectorsize = 1; +#endif gangdist = (int *) __builtin_alloca (gangsize * sizeof (int)); workerdist = (int *) __builtin_alloca (workersize * sizeof (int)); @@ -92,6 +107,11 @@ int main () int w = (ary[ix] >> 8) & 255; int v = ary[ix] & 255; + if (g >= gangsize + || w >= workersize + || v >= vectorsize) + __builtin_abort (); + gangdist[g]++; workerdist[w]++; vectordist[v]++; diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-gwv-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-gwv-1.c index c1a2d0cffe1..4099d6072da 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-gwv-1.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-gwv-1.c @@ -16,8 +16,11 @@ int main () int t = 0, h = 0; int gangsize, workersize, vectorsize; -#pragma acc parallel num_gangs(32) num_workers(32) vector_length(32) \ - copy(ondev) copyout(gangsize, workersize, vectorsize) +#define NG 32 +#define NW 32 +#define VL 32 +#pragma acc parallel num_gangs(NG) num_workers(NW) vector_length(VL) \ + copy(ondev) /* { dg-note {variable 'ix' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-2 } */ { #pragma acc loop gang worker vector reduction(+:t) @@ -42,10 +45,19 @@ int main () } t += val; } - gangsize = __builtin_goacc_parlevel_size (GOMP_DIM_GANG); - workersize = __builtin_goacc_parlevel_size (GOMP_DIM_WORKER); - vectorsize = __builtin_goacc_parlevel_size (GOMP_DIM_VECTOR); } + gangsize = NG; + workersize = NW; + vectorsize = VL; +#ifdef ACC_DEVICE_TYPE_radeon + /* AMD GCN has an upper limit of 'num_workers(16)'. */ + if (workersize > 16) + workersize = 16; + /* AMD GCN uses the autovectorizer for the vector dimension: the use + of a function call in vector-partitioned code in this test is not + currently supported. */ + vectorsize = 1; +#endif for (ix = 0; ix < N; ix++) { diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-v-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-v-1.c index 58c7b6ab57f..0fe368623c3 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-v-1.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-v-1.c @@ -17,7 +17,8 @@ int main () int t = 0, h = 0; int vectorsize; -#pragma acc parallel vector_length(32) copy(ondev) copyout(vectorsize) +#define VL 32 +#pragma acc parallel vector_length(VL) copy(ondev) /* { dg-note {variable 'ix' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-1 } */ { #pragma acc loop vector reduction (+:t) @@ -42,8 +43,14 @@ int main () } t += val; } - vectorsize = __builtin_goacc_parlevel_size (GOMP_DIM_VECTOR); } + vectorsize = VL; +#ifdef ACC_DEVICE_TYPE_radeon + /* AMD GCN uses the autovectorizer for the vector dimension: the use + of a function call in vector-partitioned code in this test is not + currently supported. */ + vectorsize = 1; +#endif for (ix = 0; ix < N; ix++) { diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-v-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-v-2.c index 85931f5e433..0cf2d473eb8 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-v-2.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-v-2.c @@ -17,7 +17,8 @@ int main () int q = 0, h = 0; int vectorsize; -#pragma acc parallel vector_length(32) copy(q) copy(ondev) copyout(vectorsize) +#define VL 32 +#pragma acc parallel vector_length(VL) copy(q) copy(ondev) /* { dg-note {variable 't' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-1 } */ /* { dg-note {variable 'ix' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-2 } */ { @@ -46,8 +47,14 @@ int main () t += val; } q = t; - vectorsize = __builtin_goacc_parlevel_size (GOMP_DIM_VECTOR); } + vectorsize = VL; +#ifdef ACC_DEVICE_TYPE_radeon + /* AMD GCN uses the autovectorizer for the vector dimension: the use + of a function call in vector-partitioned code in this test is not + currently supported. */ + vectorsize = 1; +#endif for (ix = 0; ix < N; ix++) { diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-w-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-w-1.c index b9ceec9887d..b0cb09399ba 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-w-1.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-w-1.c @@ -19,8 +19,10 @@ int main () int t = 0, h = 0; int workersize; -#pragma acc parallel num_workers(32) vector_length(32) copy(ondev) \ - copyout(workersize) +#define NW 32 +#define VL 32 +#pragma acc parallel num_workers(NW) vector_length(VL) \ + copy(ondev) /* { dg-note {variable 'ix' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-2 } */ /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .-3 } */ { @@ -46,8 +48,13 @@ int main () } t += val; } - workersize = __builtin_goacc_parlevel_size (GOMP_DIM_WORKER); } + workersize = NW; +#ifdef ACC_DEVICE_TYPE_radeon + /* AMD GCN has an upper limit of 'num_workers(16)'. */ + if (workersize > 16) + workersize = 16; +#endif for (ix = 0; ix < N; ix++) { diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-w-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-w-2.c index ff5e4a1656b..f9baedb0c46 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-w-2.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-w-2.c @@ -19,8 +19,10 @@ int main () int q = 0, h = 0; int workersize; -#pragma acc parallel num_workers(32) vector_length(32) copy(q) copy(ondev) \ - copyout(workersize) +#define NW 32 +#define VL 32 +#pragma acc parallel num_workers(NW) vector_length(VL) \ + copy(q) copy(ondev) /* { dg-note {variable 't' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-2 } */ /* { dg-note {variable 'ix' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-3 } */ /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .-4 } */ @@ -50,8 +52,13 @@ int main () t += val; } q = t; - workersize = __builtin_goacc_parlevel_size (GOMP_DIM_WORKER); } + workersize = NW; +#ifdef ACC_DEVICE_TYPE_radeon + /* AMD GCN has an upper limit of 'num_workers(16)'. */ + if (workersize > 16) + workersize = 16; +#endif for (ix = 0; ix < N; ix++) { diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-wv-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-wv-1.c index 5d60899acc1..fadb2627f73 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-wv-1.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-wv-1.c @@ -16,8 +16,10 @@ int main () int t = 0, h = 0; int workersize, vectorsize; -#pragma acc parallel num_workers(32) vector_length(32) copy(ondev) \ - copyout(workersize, vectorsize) +#define NW 32 +#define VL 32 +#pragma acc parallel num_workers(NW) vector_length(VL) \ + copy(ondev) /* { dg-note {variable 'ix' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-2 } */ { #pragma acc loop worker vector reduction (+:t) @@ -42,9 +44,18 @@ int main () } t += val; } - workersize = __builtin_goacc_parlevel_size (GOMP_DIM_WORKER); - vectorsize = __builtin_goacc_parlevel_size (GOMP_DIM_VECTOR); } + workersize = NW; + vectorsize = VL; +#ifdef ACC_DEVICE_TYPE_radeon + /* AMD GCN has an upper limit of 'num_workers(16)'. */ + if (workersize > 16) + workersize = 16; + /* AMD GCN uses the autovectorizer for the vector dimension: the use + of a function call in vector-partitioned code in this test is not + currently supported. */ + vectorsize = 1; +#endif for (ix = 0; ix < N; ix++) { diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-v-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-v-1.c index 9ccc1a89b13..bfd22b0db7f 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-v-1.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-v-1.c @@ -20,8 +20,9 @@ int main () for (ix = 0; ix < N;ix++) ary[ix] = -1; -#pragma acc parallel vector_length(32) copy(ary) copy(ondev) \ - copyout(vectorsize) +#define VL 32 +#pragma acc parallel vector_length(VL) \ + copy(ary) copy(ondev) /* { dg-note {variable 'ix' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-2 } */ { #pragma acc loop vector @@ -44,8 +45,14 @@ int main () else ary[ix] = ix; } - vectorsize = __builtin_goacc_parlevel_size (GOMP_DIM_VECTOR); } + vectorsize = VL; +#ifdef ACC_DEVICE_TYPE_radeon + /* AMD GCN uses the autovectorizer for the vector dimension: the use + of a function call in vector-partitioned code in this test is not + currently supported. */ + vectorsize = 1; +#endif for (ix = 0; ix < N; ix++) { diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-w-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-w-1.c index 0e99ec62038..b910e251fb1 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-w-1.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-w-1.c @@ -23,8 +23,10 @@ int main () for (ix = 0; ix < N;ix++) ary[ix] = -1; -#pragma acc parallel num_workers(32) vector_length(32) copy(ary) copy(ondev) \ - copyout(workersize) +#define NW 32 +#define VL 32 +#pragma acc parallel num_workers(NW) vector_length(VL) \ + copy(ary) copy(ondev) /* { dg-note {variable 'ix' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-2 } */ /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } .-3 } */ { @@ -48,8 +50,13 @@ int main () else ary[ix] = ix; } - workersize = __builtin_goacc_parlevel_size (GOMP_DIM_WORKER); } + workersize = NW; +#ifdef ACC_DEVICE_TYPE_radeon + /* AMD GCN has an upper limit of 'num_workers(16)'. */ + if (workersize > 16) + workersize = 16; +#endif for (ix = 0; ix < N; ix++) { diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-wv-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-wv-1.c index f4707d15394..77326068a66 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-wv-1.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-wv-1.c @@ -20,8 +20,10 @@ int main () for (ix = 0; ix < N;ix++) ary[ix] = -1; -#pragma acc parallel num_workers(32) vector_length(32) copy(ary) copy(ondev) \ - copyout(workersize, vectorsize) +#define NW 32 +#define VL 32 +#pragma acc parallel num_workers(NW) vector_length(VL) \ + copy(ary) copy(ondev) /* { dg-note {variable 'ix' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-2 } */ { #pragma acc loop worker vector @@ -44,9 +46,18 @@ int main () else ary[ix] = ix; } - workersize = __builtin_goacc_parlevel_size (GOMP_DIM_WORKER); - vectorsize = __builtin_goacc_parlevel_size (GOMP_DIM_VECTOR); } + workersize = NW; + vectorsize = VL; +#ifdef ACC_DEVICE_TYPE_radeon + /* AMD GCN has an upper limit of 'num_workers(16)'. */ + if (workersize > 16) + workersize = 16; + /* AMD GCN uses the autovectorizer for the vector dimension: the use + of a function call in vector-partitioned code in this test is not + currently supported. */ + vectorsize = 1; +#endif for (ix = 0; ix < N; ix++) { diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-gwv-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-gwv-1.c index da13d84908a..81e08119214 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-gwv-1.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-gwv-1.c @@ -35,14 +35,27 @@ int main () for (ix = 0; ix < N;ix++) ary[ix] = -1; -#pragma acc parallel num_gangs(32) num_workers(32) vector_length(32) copy(ary) copy(ondev) copyout(gangsize, workersize, vectorsize) +#define NG 32 +#define NW 32 +#define VL 32 +#pragma acc parallel num_gangs(NG) num_workers(NW) vector_length(VL) \ + copy(ary) copy(ondev) { ondev = acc_on_device (acc_device_not_host); gang (ary); - gangsize = __builtin_goacc_parlevel_size (GOMP_DIM_GANG); - workersize = __builtin_goacc_parlevel_size (GOMP_DIM_WORKER); - vectorsize = __builtin_goacc_parlevel_size (GOMP_DIM_VECTOR); } + gangsize = NG; + workersize = NW; + vectorsize = VL; +#ifdef ACC_DEVICE_TYPE_radeon + /* AMD GCN has an upper limit of 'num_workers(16)'. */ + if (workersize > 16) + workersize = 16; + /* AMD GCN uses the autovectorizer for the vector dimension: the use + of a function call in vector-partitioned code in this test is not + currently supported. */ + vectorsize = 1; +#endif for (ix = 0; ix < N; ix++) { diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-v-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-v-1.c index dd7bb6cdcd1..7310906bd2d 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-v-1.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-v-1.c @@ -35,13 +35,20 @@ int main () for (ix = 0; ix < N;ix++) ary[ix] = -1; -#pragma acc parallel vector_length(32) copy(ary) copy(ondev) \ - copyout(vectorsize) +#define VL 32 +#pragma acc parallel vector_length(VL) \ + copy(ary) copy(ondev) { ondev = acc_on_device (acc_device_not_host); vector (ary); - vectorsize = __builtin_goacc_parlevel_size (GOMP_DIM_VECTOR); } + vectorsize = VL; +#ifdef ACC_DEVICE_TYPE_radeon + /* AMD GCN uses the autovectorizer for the vector dimension: the use + of a function call in vector-partitioned code in this test is not + currently supported. */ + vectorsize = 1; +#endif for (ix = 0; ix < N; ix++) { diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-w-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-w-1.c index b9137d81935..4521cb91143 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-w-1.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-w-1.c @@ -39,13 +39,20 @@ int main () for (ix = 0; ix < N;ix++) ary[ix] = -1; -#pragma acc parallel num_workers(32) vector_length(32) copy(ary) copy(ondev) \ - copyout(workersize) +#define NW 32 +#define VL 32 +#pragma acc parallel num_workers(NW) vector_length(VL) \ + copy(ary) copy(ondev) { ondev = acc_on_device (acc_device_not_host); worker (ary); - workersize = __builtin_goacc_parlevel_size (GOMP_DIM_WORKER); } + workersize = NW; +#ifdef ACC_DEVICE_TYPE_radeon + /* AMD GCN has an upper limit of 'num_workers(16)'. */ + if (workersize > 16) + workersize = 16; +#endif for (ix = 0; ix < N; ix++) { diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-wv-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-wv-1.c index 73696e4e59a..647d075bb00 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-wv-1.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-wv-1.c @@ -35,14 +35,25 @@ int main () for (ix = 0; ix < N;ix++) ary[ix] = -1; -#pragma acc parallel num_workers(32) vector_length(32) copy(ary) copy(ondev) \ - copyout(workersize, vectorsize) +#define NW 32 +#define VL 32 +#pragma acc parallel num_workers(NW) vector_length(VL) \ + copy(ary) copy(ondev) { ondev = acc_on_device (acc_device_not_host); worker (ary); - workersize = __builtin_goacc_parlevel_size (GOMP_DIM_WORKER); - vectorsize = __builtin_goacc_parlevel_size (GOMP_DIM_VECTOR); } + workersize = NW; + vectorsize = VL; +#ifdef ACC_DEVICE_TYPE_radeon + /* AMD GCN has an upper limit of 'num_workers(16)'. */ + if (workersize > 16) + workersize = 16; + /* AMD GCN uses the autovectorizer for the vector dimension: the use + of a function call in vector-partitioned code in this test is not + currently supported. */ + vectorsize = 1; +#endif for (ix = 0; ix < N; ix++) { -- 2.34.1