From patchwork Wed Oct 9 11:35:12 2019 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 8bit X-Patchwork-Submitter: Thomas Schwinge X-Patchwork-Id: 1173764 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@bilbo.ozlabs.org Authentication-Results: ozlabs.org; spf=pass (mailfrom) smtp.mailfrom=gcc.gnu.org (client-ip=209.132.180.131; helo=sourceware.org; envelope-from=gcc-patches-return-510535-incoming=patchwork.ozlabs.org@gcc.gnu.org; receiver=) Authentication-Results: ozlabs.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: ozlabs.org; dkim=pass (1024-bit key; unprotected) header.d=gcc.gnu.org header.i=@gcc.gnu.org header.b="Z8gzUsWU"; dkim-atps=neutral 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 46pBwv1VF1z9s7T for ; Wed, 9 Oct 2019 22:35:40 +1100 (AEDT) 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:date:message-id:mime-version:content-type; q=dns; s=default; b=x/gaIncK6eksOaVeMisS8PHXnlXeKfnUHYZH5NnHCE+I101K0Y vkhdaJJhesp/Iqjjye6Ge8cnFpUJv19KjU5Yfe7S/tfzpTfqSpi8ib+jmMfdhBRK RDjU6GhowqT3n+B+c4l8Lf3XwfPBRIodZTSOgBzqxJRnSV3x3waqs07aY= 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:date:message-id:mime-version:content-type; s= default; bh=X0+09exCj0kFs4Pe4R5NPc9c3j8=; b=Z8gzUsWUdkVHcFPUuNZC 3rBpbqeDrknG/VM8++Wxo2tu76RTIVYrba2BG9/O74Q8YwIIe6GgWndflVLppGYm QPPPVUoJsGItNWvo+sxrLsXBM8RP49vFpIGE+6a1qZuz2B690otUn/1MtXJiY8dl AmmD8JE4MgBIljdF2iO84Wk= Received: (qmail 24885 invoked by alias); 9 Oct 2019 11:35:32 -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 24780 invoked by uid 89); 9 Oct 2019 11:35:32 -0000 Authentication-Results: sourceware.org; auth=none X-Spam-SWARE-Status: No, score=-15.9 required=5.0 tests=AWL, BAYES_00, GIT_PATCH_0, GIT_PATCH_1, GIT_PATCH_2, GIT_PATCH_3, SPF_PASS autolearn=ham version=3.3.1 spammy= X-HELO: esa3.mentor.iphmx.com Received: from esa3.mentor.iphmx.com (HELO esa3.mentor.iphmx.com) (68.232.137.180) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with ESMTP; Wed, 09 Oct 2019 11:35:31 +0000 IronPort-SDR: rX1c0qArwstP6S2RDiBG6Y3fU2RzXqN7PueiUkgqvCsuzkCc5xGBY+aXmWhXftrLIntdugVRKU s/9k0TroN4SLehbySQ0G+JodGgYphLmthRLm/qwm/A/fMVnPsdwZAXkSFcOpv0e+XCF95okaQ0 Y3XDLiYkho8Sj/pTIXXXOyouCirnsRPy7Z+ebUi3TDgdELs26FgPxPtouzAqg5I122Cq9Kd+w4 znvuKU4ZNTQulOkXOrFXhCVFZaaWyd1osVyxuAkiaKJMk/zYiMI04RRL0pHmMUhXvNhAB0TN9c 6Zk= Received: from orw-gwy-01-in.mentorg.com ([192.94.38.165]) by esa3.mentor.iphmx.com with ESMTP; 09 Oct 2019 03:35:29 -0800 IronPort-SDR: v9cosRdQT1+KVAL1f4WAJBFO6gh9SswdUV5EAzVZXJZa90eb9yX8R1zUHx+TRwrCDBkEYVttlH N2Ocm54sy96lhLv2PfbYKRlut5wqk5jysJaqDiiudWzgPTnK13qI2vHhNlWc/ICGqCmCdq8ohH N2frTW8sfpsYxPnB89wDe/TT+c8BVkPuhIY1HYxYR2zQGSlKljdUDs2mrEvr+kRAqa+5RvhCPu hc2bRNUBcqorlbNbIc+eNVyZzhnCHZvFFshRdwFeNu9iBzIbqh8SimszRzdweukJR0z2NuFPJU N3I= From: Thomas Schwinge To: CC: Julian Brown , Jakub Jelinek Subject: PR92036 "OpenACC 'firstprivate' clause: initial value" User-Agent: Notmuch/0.29.1+93~g67ed7df (https://notmuchmail.org) Emacs/26.1 (x86_64-pc-linux-gnu) Date: Wed, 9 Oct 2019 13:35:12 +0200 Message-ID: <87o8yqywwf.fsf@euler.schwinge.homeip.net> MIME-Version: 1.0 Hi! In r276757 "[PR92036] Add 'libgomp.oacc-c-c++-common/data-firstprivate-1.c'", I committed the attached to document the status quo, and then what I think it instead should be per PR92036 "OpenACC 'firstprivate' clause: initial value". Grüße Thomas From 5fc105387fda327bf5e36ae6f997e415da6d1f37 Mon Sep 17 00:00:00 2001 From: tschwinge Date: Wed, 9 Oct 2019 11:31:14 +0000 Subject: [PATCH] [PR92036] Add 'libgomp.oacc-c-c++-common/data-firstprivate-1.c' libgomp/ PR middle-end/92036 * testsuite/libgomp.oacc-c-c++-common/data-firstprivate-1.c: New file. git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/trunk@276757 138bc75d-0d04-0410-961f-82ee72b054a4 --- libgomp/ChangeLog | 6 + .../data-firstprivate-1.c | 165 ++++++++++++++++++ 2 files changed, 171 insertions(+) create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/data-firstprivate-1.c diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog index 1b43c456741..319a1911882 100644 --- a/libgomp/ChangeLog +++ b/libgomp/ChangeLog @@ -1,3 +1,9 @@ +2019-10-09 Thomas Schwinge + + PR middle-end/92036 + * testsuite/libgomp.oacc-c-c++-common/data-firstprivate-1.c: New + file. + 2019-10-09 Tobias Burnus PR testsuite/91884 diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/data-firstprivate-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/data-firstprivate-1.c new file mode 100644 index 00000000000..8900a4e070d --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/data-firstprivate-1.c @@ -0,0 +1,165 @@ +/* Test behavior of 'firstprivate' lexically vs. dynamically nested inside a + 'data' region. */ + +#include + + +#define VERIFY(x) \ + do { \ + if (!(x)) \ + abort (); \ + } while (0); + + +/* This is basically and extended version of 't2' from 'firstprivate-1.c'. */ + +int lexically_nested_val = 2; + +static void +lexically_nested () +{ +#pragma acc data \ + copy (lexically_nested_val) + { + VERIFY (lexically_nested_val == 2); + +#pragma acc parallel \ + present (lexically_nested_val) + { + VERIFY (lexically_nested_val == 2); + + /* This updates the device copy, or shared variable. */ + lexically_nested_val = 7; + } + +#if ACC_MEM_SHARED + VERIFY (lexically_nested_val == 7); +#else + VERIFY (lexically_nested_val == 2); +#endif + + /* This only updates the local/shared variable, but not the device + copy. */ + lexically_nested_val = 5; + +#pragma acc parallel \ + firstprivate (lexically_nested_val) + { +#if 1 /* Current behavior. */ + /* The 'firstprivate' copy is initialized from the device copy, or + shared variable. */ +# if ACC_MEM_SHARED + VERIFY (lexically_nested_val == 5); +# else + VERIFY (lexically_nested_val == 7); +# endif +#else /* Expected behavior per PR92036. */ + /* The 'firstprivate' copy is initialized from the local thread. */ + VERIFY (lexically_nested_val == 5); +#endif + + /* This updates the 'firstprivate' copy only, but not the shared + variable. */ + lexically_nested_val = 9; + } + + VERIFY (lexically_nested_val == 5); + } + /* If not shared, the device copy has now been copied back. */ + +#if ACC_MEM_SHARED + VERIFY (lexically_nested_val == 5); +#else + VERIFY (lexically_nested_val == 7); +#endif +} + + +int dynamically_nested_val = 2; + +/* Same as above, but compute construct 1 broken out, so no longer lexically + nested inside 'data' region. */ + +static void +dynamically_nested_compute_1 () +{ +#pragma acc parallel \ + present (dynamically_nested_val) + { + VERIFY (dynamically_nested_val == 2); + + /* This updates the device copy, or shared variable. */ + dynamically_nested_val = 7; + } +} + +/* Same as above, but compute construct 2 broken out, so no longer lexically + nested inside 'data' region. */ + +static void +dynamically_nested_compute_2 () +{ +#pragma acc parallel \ + firstprivate (dynamically_nested_val) + { +#if 1 /* Current behavior. */ + /* The 'firstprivate' copy is initialized from the device copy, or shared + variable. */ +# if ACC_MEM_SHARED + VERIFY (dynamically_nested_val == 5); +# else + VERIFY (dynamically_nested_val == 7); +# endif +#else /* Expected behavior per PR92036. */ + /* The 'firstprivate' copy is initialized from the local thread. */ + VERIFY (dynamically_nested_val == 5); +#endif + + /* This updates the 'firstprivate' copy only, but not the shared + variable. */ + dynamically_nested_val = 9; + } +} + +static void +dynamically_nested () +{ +#pragma acc data \ + copy (dynamically_nested_val) + { + VERIFY (dynamically_nested_val == 2); + + dynamically_nested_compute_1 (); + +#if ACC_MEM_SHARED + VERIFY (dynamically_nested_val == 7); +#else + VERIFY (dynamically_nested_val == 2); +#endif + + /* This only updates the local/shared variable, but not the device + copy. */ + dynamically_nested_val = 5; + + dynamically_nested_compute_2 (); + + VERIFY (dynamically_nested_val == 5); + } + /* If not shared, the device copy has now been copied back. */ + +#if ACC_MEM_SHARED + VERIFY (dynamically_nested_val == 5); +#else + VERIFY (dynamically_nested_val == 7); +#endif +} + + +int +main() +{ + lexically_nested (); + dynamically_nested (); + + return 0; +} -- 2.17.1