diff mbox series

[committed] Address compiler diagnostics in libgomp.oacc-c-c++-common/pr87835.c (was: [committed][nvptx, libgomp] Fix map_push)

Message ID 875zqlfehs.fsf@euler.schwinge.homeip.net
State New
Headers show
Series [committed] Address compiler diagnostics in libgomp.oacc-c-c++-common/pr87835.c (was: [committed][nvptx, libgomp] Fix map_push) | expand

Commit Message

Thomas Schwinge May 8, 2019, 10:05 a.m. UTC 
Hi!

On Wed, 23 Jan 2019 09:19:33 +0100, Tom de Vries <tdevries@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
diff mbox series

Patch

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  <thomas@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