From patchwork Fri Dec 14 21:18:03 2018 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 8bit X-Patchwork-Submitter: Thomas Schwinge X-Patchwork-Id: 1013715 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-492530-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="Cxq3FWLP"; 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 43Gk1G6pxjz9s3Z for ; Sat, 15 Dec 2018 08:18:26 +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 :content-transfer-encoding; q=dns; s=default; b=rlR6buplfIPMLEMd B/V79p4zGCabsT6Jtr2uE7xljZa1koThhWqoJ33k2Indlez8GGxaSySbEyD+GpdM hzPmWgJj10FSygNMJzYmzYuhARRYJc1NThv1FQ9I1tTO1t3PA0tuvD7vrV3Vamal RpB4wTKxXqF6A0lG9YyTe3CD7RA= 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 :content-transfer-encoding; s=default; bh=llhL+hBddqt89D9+FoMMjt MqSss=; b=Cxq3FWLPmw3fsYRrZLABC26EHJAJ2waEKkQVQyHK71N8zsLZlyH0de 9ZoprKjyNJ+3W4mo8oBHqyQMHL/V5+GCssiB4VmE7s4FYTHkDhqyGyPgiJ323QKc C6nyTSu5z/U2LiRQH88qkBB5Njq3WtTK9cA6EdP2Jy4IyLNFTpXXg= Received: (qmail 101264 invoked by alias); 14 Dec 2018 21:18:18 -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 101252 invoked by uid 89); 14 Dec 2018 21:18:18 -0000 Authentication-Results: sourceware.org; auth=none X-Spam-SWARE-Status: No, score=-26.9 required=5.0 tests=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.2 spammy= 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; Fri, 14 Dec 2018 21:18:15 +0000 Received: from svr-orw-mbx-01.mgc.mentorg.com ([147.34.90.201]) by relay1.mentorg.com with esmtps (TLSv1.2:ECDHE-RSA-AES256-SHA384:256) id 1gXuqc-00064y-37 from Thomas_Schwinge@mentor.com for gcc-patches@gcc.gnu.org; Fri, 14 Dec 2018 13:18:14 -0800 Received: from svr-orw-mbx-02.mgc.mentorg.com (147.34.90.202) by svr-orw-mbx-01.mgc.mentorg.com (147.34.90.201) with Microsoft SMTP Server (TLS) id 15.0.1320.4; Fri, 14 Dec 2018 13:18:11 -0800 Received: from tftp-cs (147.34.91.1) by svr-orw-mbx-02.mgc.mentorg.com (147.34.90.202) with Microsoft SMTP Server id 15.0.1320.4 via Frontend Transport; Fri, 14 Dec 2018 13:18:11 -0800 Received: by tftp-cs (Postfix, from userid 49978) id 4064AC2321; Fri, 14 Dec 2018 13:18:11 -0800 (PST) From: Thomas Schwinge To: CC: Chung-Lin Tang Subject: [PR88495] An OpenACC async queue is always synchronized with itself User-Agent: Notmuch/0.9-125-g4686d11 (http://notmuchmail.org) Emacs/25.2.2 (x86_64-pc-linux-gnu) Date: Fri, 14 Dec 2018 22:18:03 +0100 Message-ID: <87h8ffreno.fsf@euler.schwinge.homeip.net> MIME-Version: 1.0 Hi! Committed to trunk in r267152: commit 963e7a8d58a248f8093947e9a5ba56306d36a8e2 Author: tschwinge Date: Fri Dec 14 20:43:02 2018 +0000 [PR88495] An OpenACC async queue is always synchronized with itself An OpenACC async queue is always synchronized with itself, so invocations like "#pragma acc wait(0) async(0)", or "acc_wait_async (0, 0)" don't make a lot of sense, but are still valid. libgomp/ PR libgomp/88495 * plugin/plugin-nvptx.c (nvptx_wait_async): Don't refuse "identical parameters". * testsuite/libgomp.oacc-c-c++-common/asyncwait-nop-1.c: Update. * testsuite/libgomp.oacc-c-c++-common/lib-80.c: Remove. git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/trunk@267152 138bc75d-0d04-0410-961f-82ee72b054a4 --- libgomp/ChangeLog | 6 + libgomp/plugin/plugin-nvptx.c | 3 +- .../libgomp.oacc-c-c++-common/asyncwait-nop-1.c | 3 - .../testsuite/libgomp.oacc-c-c++-common/lib-80.c | 135 --------------------- 4 files changed, 8 insertions(+), 139 deletions(-) Grüße Thomas diff --git libgomp/ChangeLog libgomp/ChangeLog index 2914066f7532..349497d58ee6 100644 --- libgomp/ChangeLog +++ libgomp/ChangeLog @@ -1,5 +1,11 @@ 2018-12-14 Thomas Schwinge + PR libgomp/88495 + * plugin/plugin-nvptx.c (nvptx_wait_async): Don't refuse + "identical parameters". + * testsuite/libgomp.oacc-c-c++-common/asyncwait-nop-1.c: Update. + * testsuite/libgomp.oacc-c-c++-common/lib-80.c: Remove. + PR libgomp/88484 * oacc-parallel.c (GOACC_wait): Correct handling for "async >= 0". * testsuite/libgomp.oacc-c-c++-common/asyncwait-nop-1.c: New file. diff --git libgomp/plugin/plugin-nvptx.c libgomp/plugin/plugin-nvptx.c index 6f9b16634b10..fb686de73f25 100644 --- libgomp/plugin/plugin-nvptx.c +++ libgomp/plugin/plugin-nvptx.c @@ -1617,8 +1617,9 @@ nvptx_wait_async (int async1, int async2) necessarily have to exist already. */ s2 = select_stream_for_async (async2, self, true, NULL); + /* A stream is always synchronized with itself. */ if (s1 == s2) - GOMP_PLUGIN_fatal ("identical parameters"); + return; e = (CUevent *) GOMP_PLUGIN_malloc (sizeof (CUevent)); diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/asyncwait-nop-1.c libgomp/testsuite/libgomp.oacc-c-c++-common/asyncwait-nop-1.c index e4f627d38bc2..4ab67363ba67 100644 --- libgomp/testsuite/libgomp.oacc-c-c++-common/asyncwait-nop-1.c +++ libgomp/testsuite/libgomp.oacc-c-c++-common/asyncwait-nop-1.c @@ -51,9 +51,6 @@ main () { for (size_t j = 0; j < values_n; ++j) { - if (values[i] == values[j]) - continue; - #pragma acc parallel wait (values[i]) async (values[j]) ; #pragma acc wait (values[i]) async (values[j]) diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/lib-80.c libgomp/testsuite/libgomp.oacc-c-c++-common/lib-80.c deleted file mode 100644 index 9a9a837fa4f2..000000000000 --- libgomp/testsuite/libgomp.oacc-c-c++-common/lib-80.c +++ /dev/null @@ -1,135 +0,0 @@ -/* { dg-do run { target openacc_nvidia_accel_selected } } */ -/* { dg-additional-options "-lcuda" } */ - -#include -#include -#include -#include -#include -#include "timer.h" - -int -main (int argc, char **argv) -{ - CUdevice dev; - CUfunction delay; - CUmodule module; - CUresult r; - CUstream stream; - int N; - int i; - unsigned long *a, *d_a, dticks; - int nbytes; - float atime, dtime; - void *kargs[2]; - int clkrate; - int devnum, nprocs; - - acc_init (acc_device_nvidia); - - devnum = acc_get_device_num (acc_device_nvidia); - - r = cuDeviceGet (&dev, devnum); - if (r != CUDA_SUCCESS) - { - fprintf (stderr, "cuDeviceGet failed: %d\n", r); - abort (); - } - - r = - cuDeviceGetAttribute (&nprocs, CU_DEVICE_ATTRIBUTE_MULTIPROCESSOR_COUNT, - dev); - if (r != CUDA_SUCCESS) - { - fprintf (stderr, "cuDeviceGetAttribute failed: %d\n", r); - abort (); - } - - r = cuDeviceGetAttribute (&clkrate, CU_DEVICE_ATTRIBUTE_CLOCK_RATE, dev); - if (r != CUDA_SUCCESS) - { - fprintf (stderr, "cuDeviceGetAttribute failed: %d\n", r); - abort (); - } - - r = cuModuleLoad (&module, "subr.ptx"); - if (r != CUDA_SUCCESS) - { - fprintf (stderr, "cuModuleLoad failed: %d\n", r); - abort (); - } - - r = cuModuleGetFunction (&delay, module, "delay"); - if (r != CUDA_SUCCESS) - { - fprintf (stderr, "cuModuleGetFunction failed: %d\n", r); - abort (); - } - - nbytes = nprocs * sizeof (unsigned long); - - dtime = 200.0; - - dticks = (unsigned long) (dtime * clkrate); - - N = nprocs; - - a = (unsigned long *) malloc (nbytes); - d_a = (unsigned long *) acc_malloc (nbytes); - - acc_map_data (a, d_a, nbytes); - - r = cuStreamCreate (&stream, CU_STREAM_DEFAULT); - if (r != CUDA_SUCCESS) - { - fprintf (stderr, "cuStreamCreate failed: %d\n", r); - abort (); - } - - acc_set_cuda_stream (1, stream); - - init_timers (1); - - kargs[0] = (void *) &d_a; - kargs[1] = (void *) &dticks; - - start_timer (0); - - for (i = 0; i < N; i++) - { - r = cuLaunchKernel (delay, 1, 1, 1, 1, 1, 1, 0, stream, kargs, 0); - if (r != CUDA_SUCCESS) - { - fprintf (stderr, "cuLaunchKernel failed: %d\n", r); - abort (); - } - } - - fprintf (stderr, "CheCKpOInT\n"); - acc_wait_async (1, 1); - - acc_wait (1); - - atime = stop_timer (0); - - if (atime < dtime) - { - fprintf (stderr, "actual time < delay time\n"); - abort (); - } - - acc_unmap_data (a); - - fini_timers (); - - free (a); - acc_free (d_a); - - acc_shutdown (acc_device_nvidia); - - return 0; -} - -/* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */ -/* { dg-output "identical parameters" } */ -/* { dg-shouldfail "" } */