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 625F43947C00 for ; Tue, 2 Mar 2021 12:20:34 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.3.2 sourceware.org 625F43947C00 Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=Julian_Brown@mentor.com IronPort-SDR: jVCjev9/DU2DC+TaSOWbJy86ZxgWhHgt2spbpUuaTYnEcETiJ2/I/JWD/MDEIPG2Kupv7ZMCYA IRJclDm76YYRcCVGVJJp0rYT3laS4s4N5c5asPGO/oFI5A+9J67vTttJrp23DSKjTGxTX4vRxH bZccMTBDt+Y/Ls8CXdH5PBWx4Bth6MUVTOd2rzmXvld2J8SXXdERRzKt0CAYF5LMwIyMy4DClw nyvvyDmoLwCaSsP3x8SMy8PpwkkqpnENg9C6b7RonYxLrFxrYOL9MMxKOE60Jt68G772rYTfzs Y5s= X-IronPort-AV: E=Sophos;i="5.81,216,1610438400"; d="scan'208";a="58655171" Received: from orw-gwy-02-in.mentorg.com ([192.94.38.167]) by esa3.mentor.iphmx.com with ESMTP; 02 Mar 2021 04:20:33 -0800 IronPort-SDR: k07NswysahSsQ/sATSK37wycRukmSlzh2DWf8VKLJh709bSIxzalOw9TYaPCxvTdEpbKxBvnyy w6t3s4Y/tYRyszqm0oS+PWdqsAq6uD0JlFjtYL7neYjGITdPSL9ytUZC36l7h6F6VbO8fYnoJu i/1veU3V/SoxA78KxEZAv8kmuGVEmS08wNGD2Tl08BhnkORwaHPXnP5wnby0TQRqUgqn1LuvY1 r1irUqaV4rA88HNkdwZ0cFe/H/mAavZDvD5DwPwC9CTMIeT0LzN2ZXssnXoAn+qv9d5to3PKAl nPo= From: Julian Brown To: CC: Thomas Schwinge , Tobias Burnus , Kwok Cheung Yeung , Jakub Jelinek Subject: [PATCH 2/4] openacc: Fix async bugs in several OpenACC test cases Date: Tue, 2 Mar 2021 04:20:12 -0800 Message-ID: <85f6930efb47ac6b55aa28ff5cfa244010cc5809.1614685766.git.julian@codesourcery.com> 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-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.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, 02 Mar 2021 12:20:35 -0000 Enabling worker-partitioning support in the middle end (for AMD GCN) reveals several bugs in existing tests relating to async usage. This patch fixes those up. Tested with offloading to AMD GCN. OK for stage 1? (Or now?) Julian 2021-03-02 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