[PR88495] An OpenACC async queue is always synchronized with itself
Thomas Schwinge
thomas@codesourcery.com
Fri Dec 14 21:18:00 GMT 2018
Hi!
Committed to trunk in r267152:
commit 963e7a8d58a248f8093947e9a5ba56306d36a8e2
Author: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
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(-)
diff --git libgomp/ChangeLog libgomp/ChangeLog
index 2914066f7532..349497d58ee6 100644
--- libgomp/ChangeLog
+++ libgomp/ChangeLog
@@ -1,5 +1,11 @@
2018-12-14 Thomas Schwinge <thomas@codesourcery.com>
+ 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 <stdio.h>
-#include <stdlib.h>
-#include <unistd.h>
-#include <openacc.h>
-#include <cuda.h>
-#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 "" } */
Grüße
Thomas
More information about the Gcc-patches
mailing list