[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)
Related show

Commit Message

Thomas Schwinge May 8, 2019, 10:05 a.m.
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

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