git
/
gcc
mirror of git://gcc.gnu.org/git/gcc.git
1
0
Fork 0
Code Issues Packages Projects Releases Wiki Activity

[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.

From-SVN: r267152
This commit is contained in:
Thomas Schwinge 2018-12-14 21:43:02 +01:00 committed by Thomas Schwinge
parent c8ab8aab9f
commit f847198ec3
4 changed files with 8 additions and 139 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

6
libgomp/ChangeLog
View File

@ -1,5 +1,11 @@
2018-12-14 Thomas Schwinge <thomas@codesourcery.com> 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 PR libgomp/88484
* oacc-parallel.c (GOACC_wait): Correct handling for "async >= 0". * oacc-parallel.c (GOACC_wait): Correct handling for "async >= 0".
* testsuite/libgomp.oacc-c-c++-common/asyncwait-nop-1.c: New file. * testsuite/libgomp.oacc-c-c++-common/asyncwait-nop-1.c: New file.

3
libgomp/plugin/plugin-nvptx.c
View File

@ -1617,8 +1617,9 @@ nvptx_wait_async (int async1, int async2)
necessarily have to exist already. */ necessarily have to exist already. */
s2 = select_stream_for_async (async2, self, true, NULL); s2 = select_stream_for_async (async2, self, true, NULL);
/* A stream is always synchronized with itself. */
if (s1 == s2) if (s1 == s2)
GOMP_PLUGIN_fatal ("identical parameters"); return;
e = (CUevent *) GOMP_PLUGIN_malloc (sizeof (CUevent)); e = (CUevent *) GOMP_PLUGIN_malloc (sizeof (CUevent));

3
libgomp/testsuite/libgomp.oacc-c-c++-common/asyncwait-nop-1.c
View File

@ -51,9 +51,6 @@ main ()
{ {
for (size_t j = 0; j < values_n; ++j) 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 parallel wait (values[i]) async (values[j])
; ;
#pragma acc wait (values[i]) async (values[j]) #pragma acc wait (values[i]) async (values[j])

135
libgomp/testsuite/libgomp.oacc-c-c++-common/lib-80.c
View File

@ -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 "" } */