From patchwork Wed May 19 12:10:01 2021 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 8bit X-Patchwork-Submitter: Thomas Schwinge X-Patchwork-Id: 1480870 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@bilbo.ozlabs.org Authentication-Results: ozlabs.org; spf=pass (sender SPF authorized) smtp.mailfrom=gcc.gnu.org (client-ip=2620:52:3:1:0:246e:9693:128c; helo=sourceware.org; envelope-from=gcc-patches-bounces@gcc.gnu.org; receiver=) Received: from sourceware.org (server2.sourceware.org [IPv6:2620:52:3:1:0:246e:9693:128c]) (using TLSv1.3 with cipher TLS_AES_256_GCM_SHA384 (256/256 bits) key-exchange X25519 server-signature RSA-PSS (4096 bits) server-digest SHA256) (No client certificate requested) by ozlabs.org (Postfix) with ESMTPS id 4FlWsS5ZbQz9sRf for ; Wed, 19 May 2021 22:10:20 +1000 (AEST) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id DD9023939C17; Wed, 19 May 2021 12:10:15 +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 AEB283938C27 for ; Wed, 19 May 2021 12:10:12 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.3.2 sourceware.org AEB283938C27 Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=Thomas_Schwinge@mentor.com IronPort-SDR: ElcZL+UdoMDprYpgaWV0T0rIgshHWUSy1JlvRjrHs7u26kciuNqhLfYMiVQF68XI7Ax8ggv5UM My3H+R6Ec4TsGgkRzIETpnSXFyHPfIo/plsFwJZbNOOa54h5WDfVCTLWyRQmS+NQkU84qv5HoV 4n4p/u4U/OwSZ7fJ9towzGH/R5p5vwOjvcMwh5NiZBi29VBS4/fK3OzxVo5wHhSB85jsNDt5NR /JXpdU0QopjcPq/FSL3oXmKeZeujDl0pyqJCLkYkD5wNTC6E8nBeuY7gUJZ+omE+2m+5d0xeTO FxQ= X-IronPort-AV: E=Sophos;i="5.82,312,1613462400"; d="scan'208,223";a="61381136" Received: from orw-gwy-02-in.mentorg.com ([192.94.38.167]) by esa3.mentor.iphmx.com with ESMTP; 19 May 2021 04:10:11 -0800 IronPort-SDR: k80juiEQ+fjyQgD2Y/+lXYP+8hLNPNOOlRfEE1dNYjkHUdvThx0YmLGuZuspVsz/jtkcHUYYdh /yuBd5OagcQmKpyeFkCA3UdcICx9jES9+cUvuLOV1eMeu8hcJWP9DUY4r+J42E5ac5qgoecI/u 9RTzp7NfbHeoNmiJNO0DxveqKDAUfKnlUfBJnXCCpk/uIHMo4Wk/zsdYKMHXgnG+mdrm+zWr2E NTDu7DmtFoQh5VFL7+8og5H6L7WagtRLHWuDKn+DrpHI+0uq7EQvH5pIuPLHvNheANAsW9iwy9 luo= From: Thomas Schwinge To: , Julian Brown Subject: Add 'libgomp.oacc-c-c++-common/loop-gwv-2.c' (was: [PATCH, OpenACC] Add support for gang local storage allocation in shared memory) In-Reply-To: <20180813214150.02b08598@squid.athome> References: <20180813172151.6bfcece3@squid.athome> <524d9017-e7f0-87a2-6a62-9b23abd65ac2@codesourcery.com> <20180813214150.02b08598@squid.athome> User-Agent: Notmuch/0.29.3+94~g74c3f1b (https://notmuchmail.org) Emacs/27.1 (x86_64-pc-linux-gnu) Date: Wed, 19 May 2021 14:10:01 +0200 Message-ID: <8735uiudpi.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.5 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.2 X-Spam-Checker-Version: SpamAssassin 3.4.2 (2018-09-13) 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@gcc.gnu.org Sender: "Gcc-patches" Hi! On 2018-08-13T21:41:50+0100, Julian Brown wrote: > On Mon, 13 Aug 2018 11:42:26 -0700 Cesar Philippidis wrote: >> On 08/13/2018 09:21 AM, Julian Brown wrote: >> > 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 >> > new file mode 100644 >> > index 0000000..2fa708a >> > --- /dev/null >> > +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-gwv-2.c >> > @@ -0,0 +1,106 @@ >> > +/* { dg-xfail-run-if "gangprivate failure" { openacc_nvidia_accel_selected } { "-O0" } { "" } } */ >> is the above xfail still necessary? It seems to xpass >> for me on nvptx. However, I see this regression on the host: >> >> FAIL: libgomp.oacc-c/../libgomp.oacc-c-c++-common/loop-gwv-2.c >> -DACC_DEVICE_TYPE_host=1 -DACC_MEM_SHARED=1 -O2 execution test > Oops, this was the version of the patch I meant to post (and the one I > tested). The XFAIL on loop-gwv-2.c isn't necessary, plus that test > needed some other fixes to make it pass for NVPTX (it was written for > GCN to start with). As I should find out later, this testcase actually does work without the code changes (OpenACC privatization levels) that it's accompanying -- and I don't actually see anything in the testcase that the code changes would trigger for. Maybe it was for some earlier revision of these code changes? Anyway, as it's all-PASS for all systems that I've tested on, I've now pushed "Add 'libgomp.oacc-c-c++-common/loop-gwv-2.c'" to master branch in commit 5a16fb19e7c4274f8dd9bbdd30d7d06fe2eff8af, see attached. Grüße Thomas ----------------- Mentor Graphics (Deutschland) GmbH, Arnulfstrasse 201, 80634 München Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Frank Thürauf From 5a16fb19e7c4274f8dd9bbdd30d7d06fe2eff8af Mon Sep 17 00:00:00 2001 From: Julian Brown Date: Mon, 13 Aug 2018 21:41:50 +0100 Subject: [PATCH] Add 'libgomp.oacc-c-c++-common/loop-gwv-2.c' libgomp/ * testsuite/libgomp.oacc-c-c++-common/loop-gwv-2.c: New. --- .../libgomp.oacc-c-c++-common/loop-gwv-2.c | 95 +++++++++++++++++++ 1 file changed, 95 insertions(+) create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/loop-gwv-2.c 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 new file mode 100644 index 00000000000..a4f81a39e24 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-gwv-2.c @@ -0,0 +1,95 @@ +#include +#include +#include +#include +#include +#include + +#if 0 +#define DEBUG(DIM, IDX, VAL) \ + fprintf (stderr, "%sdist[%d] = %d\n", (DIM), (IDX), (VAL)) +#else +#define DEBUG(DIM, IDX, VAL) +#endif + +#define N (32*32*32) + +int +check (const char *dim, int *dist, int dimsize) +{ + int ix; + int exit = 0; + + for (ix = 0; ix < dimsize; ix++) + { + DEBUG(dim, ix, dist[ix]); + if (dist[ix] < (N) / (dimsize + 0.5) + || dist[ix] > (N) / (dimsize - 0.5)) + { + fprintf (stderr, "did not distribute to %ss (%d not between %d " + "and %d)\n", dim, dist[ix], (int) ((N) / (dimsize + 0.5)), + (int) ((N) / (dimsize - 0.5))); + exit |= 1; + } + } + + return exit; +} + +int main () +{ + int ary[N]; + int ix; + int exit = 0; + int gangsize = 0, workersize = 0, vectorsize = 0; + 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) + { +#pragma acc loop gang worker vector + for (unsigned ix = 0; ix < N; ix++) + { + int g, w, v; + + g = __builtin_goacc_parlevel_id (GOMP_DIM_GANG); + w = __builtin_goacc_parlevel_id (GOMP_DIM_WORKER); + v = __builtin_goacc_parlevel_id (GOMP_DIM_VECTOR); + + 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); + } + + gangdist = (int *) alloca (gangsize * sizeof (int)); + workerdist = (int *) alloca (workersize * sizeof (int)); + vectordist = (int *) alloca (vectorsize * sizeof (int)); + memset (gangdist, 0, gangsize * sizeof (int)); + memset (workerdist, 0, workersize * sizeof (int)); + memset (vectordist, 0, vectorsize * sizeof (int)); + + /* Test that work is shared approximately equally amongst each active + gang/worker/vector. */ + for (ix = 0; ix < N; ix++) + { + int g = (ary[ix] >> 16) & 255; + int w = (ary[ix] >> 8) & 255; + int v = ary[ix] & 255; + + gangdist[g]++; + workerdist[w]++; + vectordist[v]++; + } + + exit = check ("gang", gangdist, gangsize); + exit |= check ("worker", workerdist, workersize); + exit |= check ("vector", vectordist, vectorsize); + + return exit; +} -- 2.30.2