Hi! On Wed, 23 Jan 2019 09:19:33 +0100, Tom de Vries <tdevr...@suse.de> 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 <openacc.h>
> +#include <stdlib.h>
> +#include "cuda.h"
> +
> +#include <stdio.h>
> +
> +#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 253ef38b3c248b69e8ab493b19b1585f291c9843 Mon Sep 17 00:00:00 2001
From: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
Date: Wed, 8 May 2019 10:01:30 +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.
git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/trunk@271004 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 64e0a8ad8df..a8ce3c241fc 100644
--- a/libgomp/ChangeLog
+++ b/libgomp/ChangeLog
@@ -1,3 +1,8 @@
+2019-05-07 Thomas Schwinge <tho...@codesourcery.com>
+
+ PR target/87835
+ * testsuite/libgomp.oacc-c-c++-common/pr87835.c: Update.
+
2019-05-06 Thomas Schwinge <tho...@codesourcery.com>
* oacc-parallel.c: Add comments to legacy entry points (GCC 5).
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
From 9f852e24d6d75f00ccca80acb5a6804912a33282 Mon Sep 17 00:00:00 2001
From: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
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 <tho...@codesourcery.com>
+
+ 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
signature.asc
Description: PGP signature
