From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: from esa3.mentor.iphmx.com (esa3.mentor.iphmx.com [68.232.137.180]) by sourceware.org (Postfix) with ESMTPS id 5DEAA385381A for ; Tue, 29 Jun 2021 23:42:23 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.1 sourceware.org 5DEAA385381A 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: 9425yVOA6fA0gwzSOm3GHtiFXHKi3V7xOPy8o1Ry+613MAa5KzLShQdXBDSJSr06/W4kctmz3/ cdYZqF167bDVFffr7lzKSn1R03Zz6nzAyZtI80BfVM81J5iFVyscv2ktgwJyAvcJldYAm60f4x lqw/QBBFuUZOaZJ+IDALK7Mcx6oJ/BPR9Xpa5F0Gmq7qRpOUNv6Hv6IOlFFogDZQC9e3+p0XPr c7dgYmGR3hxcqTrf9kg/Wom8PVCGVI2bqvIU2Gn/3YcGWfOWbHcKdGfdHP4bHdEGXxp1hHrlLk HmM= X-IronPort-AV: E=Sophos;i="5.83,310,1616486400"; d="scan'208";a="62949880" Received: from orw-gwy-01-in.mentorg.com ([192.94.38.165]) by esa3.mentor.iphmx.com with ESMTP; 29 Jun 2021 15:42:22 -0800 IronPort-SDR: oKDz6WhGRWNRBuniZqDckZJYbZcbkQI5RLTfwjVb2XCFEO+KxFEzZD8NPBpWZ7YDJ1QVByMk9/ 6JnTJChoHFsQTVzdlxhPpgV/oGmCHYqQIRXH3otHsuMFmqKHVkCDGV6poHerwbb/99AYsOUv1s cC8RYbAX30m2Ix98niFH7if1YhdY5jRbdbZs7Paq9ymyEa/koW3Da5AQKoDiLpGTCu6yyjC4EW HXlgdF9+KGnR9GhYzQGa5pd7pBkvfcp8XYU8iMJY0mDa6Apk6qaJkNDYd+Qjz/aeRycGmV39b9 vk8= From: Julian Brown To: CC: Thomas Schwinge , Jakub Jelinek , Chung-Lin Tang Subject: [PATCH 2/4] openacc: Fix async bugs in several OpenACC test cases Date: Tue, 29 Jun 2021 16:42:02 -0700 Message-ID: X-Mailer: git-send-email 2.29.2 In-Reply-To: References: MIME-Version: 1.0 Content-Transfer-Encoding: 8bit Content-Type: text/plain X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: SVR-IES-MBX-04.mgc.mentorg.com (139.181.222.4) To SVR-IES-MBX-04.mgc.mentorg.com (139.181.222.4) X-Spam-Status: No, score=-11.9 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: , X-List-Received-Date: Tue, 29 Jun 2021 23:42:24 -0000 Several OpenACC tests accidentally abuse async semantics, leading to race conditions & test failures. This patch fixes those tests. Tested with offloading to AMD GCN. I can probably self-approve this as a testcase change only, unless anyone objects. Thanks, Julian 2021-06-29 Julian Brown libgomp/ * testsuite/libgomp.oacc-c-c++-common/deep-copy-10.c: Fix async behaviour and increase number of iterations. * testsuite/libgomp.oacc-fortran/lib-16-2.f90: Fix async behaviour. * testsuite/libgomp.oacc-fortran/lib-16.f90: Likewise. --- .../libgomp.oacc-c-c++-common/deep-copy-10.c | 14 ++++++++------ .../testsuite/libgomp.oacc-fortran/lib-16-2.f90 | 5 +++++ libgomp/testsuite/libgomp.oacc-fortran/lib-16.f90 | 5 +++++ 3 files changed, 18 insertions(+), 6 deletions(-) diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-10.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-10.c index 573a8214bf0..dadb6d37942 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-10.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-10.c @@ -1,6 +1,8 @@ #include -/* Test asyncronous attach and detach operation. */ +#define ITERATIONS 1023 + +/* Test asynchronous attach and detach operation. */ typedef struct { int *a; @@ -25,13 +27,13 @@ main (int argc, char* argv[]) #pragma acc enter data copyin(m) - for (int i = 0; i < 99; i++) + for (int i = 0; i < ITERATIONS; i++) { int j; -#pragma acc parallel loop copy(m.a[0:N]) async(i % 2) +#pragma acc parallel loop copy(m.a[0:N]) async(0) for (j = 0; j < N; j++) m.a[j]++; -#pragma acc parallel loop copy(m.b[0:N]) async((i + 1) % 2) +#pragma acc parallel loop copy(m.b[0:N]) async(1) for (j = 0; j < N; j++) m.b[j]++; } @@ -40,9 +42,9 @@ main (int argc, char* argv[]) for (i = 0; i < N; i++) { - if (m.a[i] != 99) + if (m.a[i] != ITERATIONS) abort (); - if (m.b[i] != 99) + if (m.b[i] != ITERATIONS) abort (); } diff --git a/libgomp/testsuite/libgomp.oacc-fortran/lib-16-2.f90 b/libgomp/testsuite/libgomp.oacc-fortran/lib-16-2.f90 index ddd557d3be0..e2e47c967fa 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/lib-16-2.f90 +++ b/libgomp/testsuite/libgomp.oacc-fortran/lib-16-2.f90 @@ -27,6 +27,9 @@ program main if (acc_is_present (h) .neqv. .TRUE.) stop 1 + ! We must wait for the update to be done. + call acc_wait (async) + h(:) = 0 call acc_copyout_async (h, sizeof (h), async) @@ -45,6 +48,8 @@ program main if (acc_is_present (h) .neqv. .TRUE.) stop 3 + call acc_wait (async) + do i = 1, N if (h(i) /= i + i) stop 4 end do diff --git a/libgomp/testsuite/libgomp.oacc-fortran/lib-16.f90 b/libgomp/testsuite/libgomp.oacc-fortran/lib-16.f90 index ccd1ce6ee18..ef9a6f6626c 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/lib-16.f90 +++ b/libgomp/testsuite/libgomp.oacc-fortran/lib-16.f90 @@ -27,6 +27,9 @@ program main if (acc_is_present (h) .neqv. .TRUE.) stop 1 + ! We must wait for the update to be done. + call acc_wait (async) + h(:) = 0 call acc_copyout_async (h, sizeof (h), async) @@ -45,6 +48,8 @@ program main if (acc_is_present (h) .neqv. .TRUE.) stop 3 + call acc_wait (async) + do i = 1, N if (h(i) /= i + i) stop 4 end do -- 2.29.2