From patchwork Wed May 8 10:05:19 2019 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 8bit X-Patchwork-Submitter: Thomas Schwinge X-Patchwork-Id: 1096870 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-500301-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="w/LM10tr"; 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 44zXDB4Sthz9s7h for ; Wed, 8 May 2019 20:05:44 +1000 (AEST) 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:subject:in-reply-to:references:date:message-id:mime-version :content-type; q=dns; s=default; b=yfIIF9hVZ8EOK8AbBBhH5ThlvGPSk c9KEFztPTs1P7lqeUyu1ut/bo7CMIuHkpGM/Uu1G2Qel+Ieo55lUNTZ7W8uo1d74 QyUlYa0oFelrNMu0kYfWrvkxz67CU86Ja+EkFeBFVZJH1sc4zJa/1jA603jzpeQk bcl6Liw28JpH2E= 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:subject:in-reply-to:references:date:message-id:mime-version :content-type; s=default; bh=hkyVjMJFxqhCmnZq1QEdZmkI5+U=; b=w/L M10trYoq7CSqGzUGkIFilT7o1uWIlWNJ7uRhLfh1KDtAjRy3PiSg9EJ+fP7ECZT+ qq8stxPFDDDHeysv7/SyzpEzH+KUouHFswkgcXcgG3hUcZlFc0SYAKWTke5Pp0ge lv5J48Bfh8YXoKcc2bZV3XRyL4D/uxcgcwhx8Ss0= Received: (qmail 94011 invoked by alias); 8 May 2019 10:05:37 -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 94003 invoked by uid 89); 8 May 2019 10:05:36 -0000 Authentication-Results: sourceware.org; auth=none X-Spam-SWARE-Status: No, score=-15.5 required=5.0 tests=AWL, BAYES_00, GIT_PATCH_0, GIT_PATCH_1, GIT_PATCH_2, GIT_PATCH_3, RCVD_IN_DNSWL_NONE, SPF_PASS autolearn=ham version=3.3.1 spammy=U*thomas, thomas@codesourcery.com, sk:thomas@, thomascodesourcerycom X-HELO: relay1.mentorg.com Received: from relay1.mentorg.com (HELO relay1.mentorg.com) (192.94.38.131) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with ESMTP; Wed, 08 May 2019 10:05:34 +0000 Received: from svr-orw-mbx-05.mgc.mentorg.com ([147.34.90.205]) by relay1.mentorg.com with esmtps (TLSv1.2:ECDHE-RSA-AES256-SHA384:256) id 1hOJS7-00001g-Nw from Thomas_Schwinge@mentor.com ; Wed, 08 May 2019 03:05:31 -0700 Received: from SVR-ORW-MBX-09.mgc.mentorg.com (147.34.90.209) by SVR-ORW-MBX-05.mgc.mentorg.com (147.34.90.205) with Microsoft SMTP Server (TLS) id 15.0.1320.4; Wed, 8 May 2019 03:05:29 -0700 Received: from tftp-cs (147.34.91.1) by SVR-ORW-MBX-09.mgc.mentorg.com (147.34.90.209) with Microsoft SMTP Server id 15.0.1320.4 via Frontend Transport; Wed, 8 May 2019 03:05:29 -0700 Received: by tftp-cs (Postfix, from userid 49978) id 01DCCC2212; Wed, 8 May 2019 03:05:28 -0700 (PDT) From: Thomas Schwinge To: Tom de Vries , Subject: [committed] Address compiler diagnostics in libgomp.oacc-c-c++-common/pr87835.c (was: [committed][nvptx, libgomp] Fix map_push) In-Reply-To: <20190123081931.GA7156@delia> References: <20190123081931.GA7156@delia> User-Agent: Notmuch/0.9-125-g4686d11 (http://notmuchmail.org) Emacs/26.1 (x86_64-pc-linux-gnu) Date: Wed, 8 May 2019 12:05:19 +0200 Message-ID: <875zqlfehs.fsf@euler.schwinge.homeip.net> MIME-Version: 1.0 Hi! On Wed, 23 Jan 2019 09:19:33 +0100, Tom de Vries wrote: > The map field of a struct ptx_stream is [...] > The current implemention gets at least the first and most basic scenario wrong: > [...] > This problem causes the test-case asyncwait-1.c to fail intermittently on some > systems. The pr87835.c test-case added here is a a minimized and modified > version of asyncwait-1.c (avoiding the kernel construct) that is more likely to > fail. Indeed, with one OpenACC directive fixed (see below), I've been able to reliably reproduce the failure, too, for all optimization levels I tried. > Fix this by rewriting map_pop more robustly, by: [...] Thanks, belatedly. Regarding the test case: > --- /dev/null > +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr87835.c > @@ -0,0 +1,62 @@ > +/* { dg-do run { target openacc_nvidia_accel_selected } } */ > +/* { dg-additional-options "-lcuda" } */ > + > +#include > +#include > +#include "cuda.h" > + > +#include > + > +#define n 128 > + > +int > +main (void) > +{ > + CUresult r; > + CUstream stream1; > + int N = n; > + int a[n]; > + int b[n]; source-gcc/libgomp/testsuite/libgomp.oacc-c-c++-common/pr87835.c:19:7: warning: unused variable 'b' [-Wunused-variable] 19 | int b[n]; | ^ > + int c[n]; > + > + acc_init (acc_device_nvidia); > + > + r = cuStreamCreate (&stream1, CU_STREAM_NON_BLOCKING); > + if (r != CUDA_SUCCESS) > + { > + fprintf (stderr, "cuStreamCreate failed: %d\n", r); > + abort (); > + } > + > + acc_set_cuda_stream (1, stream1); > + > + for (int i = 0; i < n; i++) > + { > + a[i] = 3; > + c[i] = 0; > + } > + > +#pragma acc data copy (a, b, c) copyin (N) > + { > +#pragma acc parallel async (1) > + ; > + > +#pragma acc parallel async (1) num_gangs (320) > + #pragma loop gang source-gcc/libgomp/testsuite/libgomp.oacc-c-c++-common/pr87835.c:45: warning: ignoring #pragma loop gang [-Wunknown-pragmas] 45 | #pragma loop gang | > + for (int ii = 0; ii < N; ii++) > + c[ii] = (a[ii] + a[N - ii - 1]); > + > +#pragma acc parallel async (1) > + #pragma acc loop seq > + for (int ii = 0; ii < n; ii++) > + a[ii] = 6; > + > +#pragma acc wait (1) > + } > + > + for (int i = 0; i < n; i++) > + if (c[i] != 6) > + abort (); > + > + return 0; > +} Addressed on trunk in r271004, and on gcc-9-branch in r271005, see attached. Grüße Thomas From 9f852e24d6d75f00ccca80acb5a6804912a33282 Mon Sep 17 00:00:00 2001 From: tschwinge Date: Wed, 8 May 2019 10:03:04 +0000 Subject: [PATCH] Address compiler diagnostics in libgomp.oacc-c-c++-common/pr87835.c source-gcc/libgomp/testsuite/libgomp.oacc-c-c++-common/pr87835.c: In function 'main': source-gcc/libgomp/testsuite/libgomp.oacc-c-c++-common/pr87835.c:45: warning: ignoring #pragma loop gang [-Wunknown-pragmas] 45 | #pragma loop gang | source-gcc/libgomp/testsuite/libgomp.oacc-c-c++-common/pr87835.c:19:7: warning: unused variable 'b' [-Wunused-variable] 19 | int b[n]; | ^ libgomp/ PR target/87835 * testsuite/libgomp.oacc-c-c++-common/pr87835.c: Update. trunk r271004 git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gcc-9-branch@271005 138bc75d-0d04-0410-961f-82ee72b054a4 --- libgomp/ChangeLog | 5 +++++ libgomp/testsuite/libgomp.oacc-c-c++-common/pr87835.c | 5 ++--- 2 files changed, 7 insertions(+), 3 deletions(-) diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog index ff585e9f742..3327856787f 100644 --- a/libgomp/ChangeLog +++ b/libgomp/ChangeLog @@ -1,3 +1,8 @@ +2019-05-07 Thomas Schwinge + + PR target/87835 + * testsuite/libgomp.oacc-c-c++-common/pr87835.c: Update. + 2019-05-03 Release Manager * GCC 9.1.0 released. diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr87835.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr87835.c index 310a485e74f..88c2c7763cc 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr87835.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr87835.c @@ -16,7 +16,6 @@ main (void) CUstream stream1; int N = n; int a[n]; - int b[n]; int c[n]; acc_init (acc_device_nvidia); @@ -36,13 +35,13 @@ main (void) c[i] = 0; } -#pragma acc data copy (a, b, c) copyin (N) +#pragma acc data copy (a, c) copyin (N) { #pragma acc parallel async (1) ; #pragma acc parallel async (1) num_gangs (320) - #pragma loop gang + #pragma acc loop gang for (int ii = 0; ii < N; ii++) c[ii] = (a[ii] + a[N - ii - 1]); -- 2.17.1