diff mbox series

[OpenACC,privatization] Reject 'static', 'external' in blocks [PR90115] (was: [Patch][OG10] omp-low.c: Avoid offload-target lto1 is-missing error by not-privatizing TREE_READONLY vars)

Message ID 87tumvyiat.fsf@euler.schwinge.homeip.net
State New
Headers show
Series [OpenACC,privatization] Reject 'static', 'external' in blocks [PR90115] (was: [Patch][OG10] omp-low.c: Avoid offload-target lto1 is-missing error by not-privatizing TREE_READONLY vars) | expand

Commit Message

Thomas Schwinge May 21, 2021, 7:54 p.m. UTC 
Hi!

Might be interesting to also research the comments at the end of the
email from an OpenMP perspective?

On 2020-07-23T17:10:54+0100, Julian Brown <julian@codesourcery.com> wrote:
> On Thu, 16 Jul 2020 15:53:54 +0200
> Tobias Burnus <tobias@codesourcery.com> wrote:
>
>> [This is an OpenACC issue but I would not completely surprised if
>> something similar could occur for OpenMP offloading as well.
>> However, the patch is for an OpenACC-specific function.]

..., and this issue only occurs in OpenACC-only handling, so shouldn't
affect OpenMP.

>> This issue occurs for libgomp.oacc-fortran/privatized-ref-2.f90, for
>> which on the device lto1 complains:
>> lto1: fatal error: /tmp/ccEGJTZN.o: section A.13.1.21 is missing
>> Here, "A.13" is a TREE_STATIC, TREE_READONLY array generated by the
>> Fortran front-end and containing the array-constructor values, i.e.
>> RHS of: array = [(-2*i, i = 1, size(array))] That testcase works on
>> the trunk or on the OG10 (= devel/omp/gcc-10) branch if one reverts
>> the patch "Re-do OpenACC private variable resolution"
>> https://gcc.gnu.org/g:2f4b477223fddb84f66e494eb88d1defbd5e04a2 which
>> is scheduled but not yet submitted for mainline inclusion.

(Now it's in,
<http://mid.mail-archive.com/87fsyfzzk4.fsf@euler.schwinge.homeip.net>
"openacc: Add support for gang local storage allocation in shared
memory", etc.)

>> The
>> offloading variable table contains the variable as "A.13.10" (which
>> works fine) and the problem-causing patch causes that the code
>> .UNIQUE (OACC_PRIVATE, 0, 0, &parm.9, &A.13); gets inserted (via the
>> then-added make_oacc_private_marker in omp-low.c). Here, the decl for
>> 'A.13' does not have a varpool_node entry – and it is not streamed
>> out as separate entity. (This IFN_ is then later processed by the
>> target lto1 via omp-offload.c's execute_oacc_device_lower – where the
>> asm_name "A.13.1" appears.)
>>
>> [While I do not completely understand why the target LTO does
>> not contain the symbol, I think the following still makes sense.
>> (I do understand why the offload var table does not contain it.)]
>>
>>
>> If the variable is TREE_READONLY, there is no need to pass it
>> through the variable-privatization bits.

(I have not yet thought in detail about 'TREE_READONLY' specifically.)

>> The current check is for VAR_P and TREE_ADDRESSABLE. For the fix,
>> one could use:
>>    !TREE_READONLY
>> or
>>    !(TREE_READONLY && TREE_STATIC)
>> or
>>    !(TREE_READONLY && (TREE_STATIC || DECL_EXTERNAL)
>> I am not sure what makes more sense. I initially used the first
>> version and then moved to the last. Thoughts?
>
> I don't know when we'd have TREE_READONLY and DECL_EXTERNAL in a
> situation where it made sense to mark the decl private using this
> mechanism -- but I'll defer to your judgement!

Such things have come up in my review of the OpenACC privatization work
('libgomp.oacc-fortran/privatized-ref-2.f90' as discussed here,
'libgomp.oacc-c-c++-common/static-variable-1.c', and a few others).

There's (a) the case of user code, and indeed code like:

    #pragma acc parallel
      {
        extern int e;
        static int s;
        [...]
      }

... are probably fine to reject.  We shall discuss that in
<https://github.com/OpenACC/openacc-spec/issues/372> "C/C++ 'static'
variables" (only visible to members of the GitHub OpenACC organization).

And then there's (b) the case of code that GCC itself synthesizes -- like
the example Tobias detailed above.  We ought to handle that "fine".
Again, I have not yet tought in detail what "fine" means here, so let's
default on doing what we've been doing before, that is, disable the new
OpenACC privatization for (b) (and simply for (a), too).

I've pushed "[OpenACC privatization] Reject 'static', 'external' in
blocks [PR90115]" to master branch in commit
325aa13996bafce0c4927876c315d1fa706d9881, see attached.


Relatedly, (c) we also need to establish and verify what's to happen for
code like:

    extern int e;
    static int s;
    [...]
    #pragma acc parallel
      {
        #pragma acc loop gang private(e, s)
        for ([...])
          [...]
      }

Is the idea that these are to be "properly privatized"?  The gang-private
instances of 'e', 's' lose the 'extern', 'static' properties?  (... and
similar for any other funny such properties that may exist?  I'm sure
Fortran has something to offer, too: 'save', etc.?)

Or, are we to get gang-private instances of 'e', 's', keeping the
'extern', 'static' properties?  ("Fun!")


Grüße
 Thomas


-----------------
Mentor Graphics (Deutschland) GmbH, Arnulfstrasse 201, 80634 München Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Frank Thürauf
diff mbox series

Patch

From 325aa13996bafce0c4927876c315d1fa706d9881 Mon Sep 17 00:00:00 2001
From: Thomas Schwinge <thomas@codesourcery.com>
Date: Fri, 21 May 2021 08:51:47 +0200
Subject: [PATCH] [OpenACC privatization] Reject 'static', 'external' in blocks
 [PR90115]

	gcc/
	PR middle-end/90115
	* omp-low.c (oacc_privatization_candidate_p): Reject 'static',
	'external' in blocks.
	gcc/testsuite/
	PR middle-end/90115
	* c-c++-common/goacc/privatization-1-compute-loop.c: Update.
	* c-c++-common/goacc/privatization-1-compute.c: Likewise.
	* c-c++-common/goacc/privatization-1-routine_gang-loop.c:
	Likewise.
	* c-c++-common/goacc/privatization-1-routine_gang.c: Likewise.
	libgomp/
	PR middle-end/90115
	* testsuite/libgomp.oacc-c-c++-common/static-variable-1.c: Update.
	* testsuite/libgomp.oacc-fortran/privatized-ref-2.f90: Likewise.
---
 gcc/omp-low.c                                 | 29 +++++++++++++++++++
 .../goacc/privatization-1-compute-loop.c      |  6 ++--
 .../goacc/privatization-1-compute.c           |  6 ++--
 .../goacc/privatization-1-routine_gang-loop.c |  6 ++--
 .../goacc/privatization-1-routine_gang.c      |  6 ++--
 .../static-variable-1.c                       | 24 +++++----------
 .../libgomp.oacc-fortran/privatized-ref-2.f90 |  4 +--
 7 files changed, 45 insertions(+), 36 deletions(-)

diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index 0d63e8243ae..e00051bd3f7 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -10192,6 +10192,9 @@  oacc_privatization_candidate_p (const location_t loc, const tree c,
 {
   dump_flags_t l_dump_flags = get_openacc_privatization_dump_flags ();
 
+  /* There is some differentiation depending on block vs. clause.  */
+  bool block = !c;
+
   bool res = true;
 
   if (res && !VAR_P (decl))
@@ -10207,6 +10210,32 @@  oacc_privatization_candidate_p (const location_t loc, const tree c,
 	}
     }
 
+  if (res && block && TREE_STATIC (decl))
+    {
+      res = false;
+
+      if (dump_enabled_p ())
+	{
+	  oacc_privatization_begin_diagnose_var (l_dump_flags, loc, c, decl);
+	  dump_printf (l_dump_flags,
+		       "isn%'t candidate for adjusting OpenACC privatization level: %s\n",
+		       "static");
+	}
+    }
+
+  if (res && block && DECL_EXTERNAL (decl))
+    {
+      res = false;
+
+      if (dump_enabled_p ())
+	{
+	  oacc_privatization_begin_diagnose_var (l_dump_flags, loc, c, decl);
+	  dump_printf (l_dump_flags,
+		       "isn%'t candidate for adjusting OpenACC privatization level: %s\n",
+		       "external");
+	}
+    }
+
   if (res && !TREE_ADDRESSABLE (decl))
     {
       res = false;
diff --git a/gcc/testsuite/c-c++-common/goacc/privatization-1-compute-loop.c b/gcc/testsuite/c-c++-common/goacc/privatization-1-compute-loop.c
index 4bfb5270690..43b39c2042f 100644
--- a/gcc/testsuite/c-c++-common/goacc/privatization-1-compute-loop.c
+++ b/gcc/testsuite/c-c++-common/goacc/privatization-1-compute-loop.c
@@ -85,10 +85,8 @@  f (int i, int j, int a)
   /* { dg-note {variable 'ss' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
   /* { dg-note {variable 'func' declared in block potentially has improper OpenACC privatization level: 'function_decl'} "TODO" { target *-*-* } l_loop$c_loop } */
   /* { dg-note {variable 'func2' declared in block potentially has improper OpenACC privatization level: 'function_decl'} "TODO" { target *-*-* } l_loop$c_loop } */
-  /* { dg-note {variable 'ext' declared in block is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_loop$c_loop }
-     { dg-note {variable 'ext' ought to be adjusted for OpenACC privatization level: 'vector'} "" { target *-*-* } l_loop$c_loop } */
-  /* { dg-note {variable 'sta' declared in block is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_loop$c_loop }
-     { dg-note {variable 'sta' ought to be adjusted for OpenACC privatization level: 'vector'} "" { target *-*-* } l_loop$c_loop } */
+  /* { dg-note {variable 'ext' declared in block isn't candidate for adjusting OpenACC privatization level: external} "" { target *-*-* } l_loop$c_loop } */
+  /* { dg-note {variable 'sta' declared in block isn't candidate for adjusting OpenACC privatization level: static} "" { target *-*-* } l_loop$c_loop } */
   /* { dg-note {variable 'xx' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
   /* { dg-note {variable 'yy' declared in block is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_loop$c_loop }
      { dg-note {variable 'yy' ought to be adjusted for OpenACC privatization level: 'vector'} "" { target *-*-* } l_loop$c_loop } */
diff --git a/gcc/testsuite/c-c++-common/goacc/privatization-1-compute.c b/gcc/testsuite/c-c++-common/goacc/privatization-1-compute.c
index 4de45e5c1ed..b7c7bff64d9 100644
--- a/gcc/testsuite/c-c++-common/goacc/privatization-1-compute.c
+++ b/gcc/testsuite/c-c++-common/goacc/privatization-1-compute.c
@@ -80,10 +80,8 @@  f (int i, int j, int a)
   /* { dg-note {variable 'ss' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */
   /* { dg-note {variable 'func' declared in block potentially has improper OpenACC privatization level: 'function_decl'} "TODO" { target *-*-* } l_compute$c_compute } */
   /* { dg-note {variable 'func2' declared in block potentially has improper OpenACC privatization level: 'function_decl'} "TODO" { target *-*-* } l_compute$c_compute } */
-  /* { dg-note {variable 'ext' declared in block is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_compute$c_compute }
-     { dg-note {variable 'ext' ought to be adjusted for OpenACC privatization level: 'gang'} "" { target *-*-* } l_compute$c_compute } */
-  /* { dg-note {variable 'sta' declared in block is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_compute$c_compute }
-     { dg-note {variable 'sta' ought to be adjusted for OpenACC privatization level: 'gang'} "" { target *-*-* } l_compute$c_compute } */
+  /* { dg-note {variable 'ext' declared in block isn't candidate for adjusting OpenACC privatization level: external} "" { target *-*-* } l_compute$c_compute } */
+  /* { dg-note {variable 'sta' declared in block isn't candidate for adjusting OpenACC privatization level: static} "" { target *-*-* } l_compute$c_compute } */
   /* { dg-note {variable 'xx' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */
   /* { dg-note {variable 'yy' declared in block is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_compute$c_compute }
      { dg-note {variable 'yy' ought to be adjusted for OpenACC privatization level: 'gang'} "" { target *-*-* } l_compute$c_compute } */
diff --git a/gcc/testsuite/c-c++-common/goacc/privatization-1-routine_gang-loop.c b/gcc/testsuite/c-c++-common/goacc/privatization-1-routine_gang-loop.c
index fcc233b0886..816e4306437 100644
--- a/gcc/testsuite/c-c++-common/goacc/privatization-1-routine_gang-loop.c
+++ b/gcc/testsuite/c-c++-common/goacc/privatization-1-routine_gang-loop.c
@@ -85,10 +85,8 @@  f (int i, int j, int a)
   /* { dg-note {variable 'ss' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
   /* { dg-note {variable 'func' declared in block potentially has improper OpenACC privatization level: 'function_decl'} "TODO" { target *-*-* } l_loop$c_loop } */
   /* { dg-note {variable 'func2' declared in block potentially has improper OpenACC privatization level: 'function_decl'} "TODO" { target *-*-* } l_loop$c_loop } */
-  /* { dg-note {variable 'ext' declared in block is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_loop$c_loop }
-     { dg-note {variable 'ext' ought to be adjusted for OpenACC privatization level: 'vector'} "" { target *-*-* } l_loop$c_loop } */
-  /* { dg-note {variable 'sta' declared in block is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_loop$c_loop }
-     { dg-note {variable 'sta' ought to be adjusted for OpenACC privatization level: 'vector'} "" { target *-*-* } l_loop$c_loop } */
+  /* { dg-note {variable 'ext' declared in block isn't candidate for adjusting OpenACC privatization level: external} "" { target *-*-* } l_loop$c_loop } */
+  /* { dg-note {variable 'sta' declared in block isn't candidate for adjusting OpenACC privatization level: static} "" { target *-*-* } l_loop$c_loop } */
   /* { dg-note {variable 'xx' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
   /* { dg-note {variable 'yy' declared in block is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_loop$c_loop }
      { dg-note {variable 'yy' ought to be adjusted for OpenACC privatization level: 'vector'} "" { target *-*-* } l_loop$c_loop } */
diff --git a/gcc/testsuite/c-c++-common/goacc/privatization-1-routine_gang.c b/gcc/testsuite/c-c++-common/goacc/privatization-1-routine_gang.c
index cd6708ff205..f9f316e4ff9 100644
--- a/gcc/testsuite/c-c++-common/goacc/privatization-1-routine_gang.c
+++ b/gcc/testsuite/c-c++-common/goacc/privatization-1-routine_gang.c
@@ -84,10 +84,8 @@  f (int i, int j, int a)
   /* { dg-note {variable 'ss' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "TODO" { xfail *-*-* } l_routine$c_routine } */
   /* { dg-note {variable 'func' declared in block potentially has improper OpenACC privatization level: 'function_decl'} "TODO" { xfail *-*-* } l_routine$c_routine } */
   /* { dg-note {variable 'func2' declared in block potentially has improper OpenACC privatization level: 'function_decl'} "TODO" { xfail *-*-* } l_routine$c_routine } */
-  /* { dg-note {variable 'ext' declared in block is candidate for adjusting OpenACC privatization level} "TODO" { xfail *-*-* } l_routine$c_routine }
-     { dg-note {variable 'ext' ought to be adjusted for OpenACC privatization level: 'gang'} "TODO" { xfail *-*-* } l_routine$c_routine } */
-  /* { dg-note {variable 'sta' declared in block is candidate for adjusting OpenACC privatization level} "TODO" { xfail *-*-* } l_routine$c_routine }
-     { dg-note {variable 'sta' ought to be adjusted for OpenACC privatization level: 'gang'} "TODO" { xfail *-*-* } l_routine$c_routine } */
+  /* { dg-note {variable 'ext' declared in block isn't candidate for adjusting OpenACC privatization level: external} "TODO" { xfail *-*-* } l_routine$c_routine } */
+  /* { dg-note {variable 'sta' declared in block isn't candidate for adjusting OpenACC privatization level: static} "TODO" { xfail *-*-* } l_routine$c_routine } */
   /* { dg-note {variable 'xx' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "TODO" { xfail *-*-* } l_routine$c_routine } */
   /* { dg-note {variable 'yy' declared in block is candidate for adjusting OpenACC privatization level} "TODO" { xfail *-*-* } l_routine$c_routine }
      { dg-note {variable 'yy' ought to be adjusted for OpenACC privatization level: 'gang'} "TODO" { xfail *-*-* } l_routine$c_routine } */
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-variable-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-variable-1.c
index 6a4c6a0e85f..ceb2c88d3e5 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-variable-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-variable-1.c
@@ -46,9 +46,7 @@  static void t0_c(void)
 #pragma acc parallel \
   reduction(max:num_gangs_actual) \
   reduction(max:result)
-      /* { dg-note {variable 'var' declared in block is candidate for adjusting OpenACC privatization level} "" { target *-*-* } .-3 }
-	 { dg-note {variable 'var' ought to be adjusted for OpenACC privatization level: 'gang'} "" { target *-*-* } .-4 }
-	 { dg-note {variable 'var' adjusted for OpenACC privatization level: 'gang'} "" { target { ! openacc_host_selected } } .-5 } */
+      /* { dg-note {variable 'var' declared in block isn't candidate for adjusting OpenACC privatization level: static} "" { target *-*-* } .-3 } */
       {
 	num_gangs_actual = 1 + __builtin_goacc_parlevel_id(GOMP_DIM_GANG);
 
@@ -143,9 +141,7 @@  static void t1_c(void)
   num_gangs(num_gangs_request) \
   reduction(max:num_gangs_actual) \
   reduction(max:result)
-      /* { dg-note {variable 'var' declared in block is candidate for adjusting OpenACC privatization level} "" { target *-*-* } .-4 }
-	 { dg-note {variable 'var' ought to be adjusted for OpenACC privatization level: 'gang'} "" { target *-*-* } .-5 }
-	 { dg-note {variable 'var' adjusted for OpenACC privatization level: 'gang'} "" { target { ! openacc_host_selected } } .-6 } */
+      /* { dg-note {variable 'var' declared in block isn't candidate for adjusting OpenACC privatization level: static} "" { target *-*-* } .-4 } */
       {
 	num_gangs_actual = 1 + __builtin_goacc_parlevel_id(GOMP_DIM_GANG);
 
@@ -317,10 +313,8 @@  static void t2(void)
   present(results_1) \
   num_gangs(num_gangs_request_1) \
   async(1)
-	/* { dg-note {variable 'var' declared in block is candidate for adjusting OpenACC privatization level} "" { target *-*-* } .-4 }
-	   { dg-note {variable 'var' ought to be adjusted for OpenACC privatization level: 'gang'} "" { target *-*-* } .-5 }
-	   { dg-note {variable 'var' adjusted for OpenACC privatization level: 'gang'} "" { target { ! openacc_host_selected } } .-6 } */
-	/* { dg-note {variable 'tmp' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-7 } */
+	/* { dg-note {variable 'var' declared in block isn't candidate for adjusting OpenACC privatization level: static} "" { target *-*-* } .-4 } */
+	/* { dg-note {variable 'tmp' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-5 } */
 	{
 	  static int var = var_init_1;
 
@@ -344,10 +338,8 @@  static void t2(void)
   present(results_3) \
   num_gangs(num_gangs_request_3) \
   async(3)
-	/* { dg-note {variable 'var' declared in block is candidate for adjusting OpenACC privatization level} "" { target *-*-* } .-4 }
-	   { dg-note {variable 'var' ought to be adjusted for OpenACC privatization level: 'gang'} "" { target *-*-* } .-5 }
-	   { dg-note {variable 'var' adjusted for OpenACC privatization level: 'gang'} "" { target { ! openacc_host_selected } } .-6 } */
-	/* { dg-note {variable 'tmp' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-7 } */
+	/* { dg-note {variable 'var' declared in block isn't candidate for adjusting OpenACC privatization level: static} "" { target *-*-* } .-4 } */
+	/* { dg-note {variable 'tmp' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-5 } */
 	{
 	  static int var = var_init_3;
 
@@ -468,9 +460,7 @@  static void pr84992_1(void)
   int n[1];
   n[0] = 3;
 #pragma acc parallel copy(n)
-  /* { dg-note {variable 'test' declared in block is candidate for adjusting OpenACC privatization level} "" { target *-*-* } .-1 }
-     { dg-note {variable 'test' ought to be adjusted for OpenACC privatization level: 'gang'} "" { target *-*-* } .-2 }
-     { dg-note {variable 'test' adjusted for OpenACC privatization level: 'gang'} "" { target { ! openacc_host_selected } } .-3 } */
+  /* { dg-note {variable 'test' declared in block isn't candidate for adjusting OpenACC privatization level: static} "" { target *-*-* } .-1 } */
   {
     static const int test[] = {1,2,3,4};
     n[0] = test[n[0]];
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/privatized-ref-2.f90 b/libgomp/testsuite/libgomp.oacc-fortran/privatized-ref-2.f90
index ba638da8628..60803e48cbe 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/privatized-ref-2.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/privatized-ref-2.f90
@@ -74,9 +74,7 @@  contains
     ! { dg-note {variable 'parm\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_compute$c_compute }
     ! { dg-note {variable 'parm\.[0-9]+' ought to be adjusted for OpenACC privatization level: 'gang'} "" { target *-*-* } l_compute$c_compute }
     ! { dg-note {variable 'parm\.[0-9]+' adjusted for OpenACC privatization level: 'gang'} "" { target { ! { openacc_host_selected || openacc_nvidia_accel_selected } } } l_compute$c_compute }
-    ! { dg-note {variable 'A\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_compute$c_compute }
-    ! { dg-note {variable 'A\.[0-9]+' ought to be adjusted for OpenACC privatization level: 'gang'} "" { target *-*-* } l_compute$c_compute }
-    ! { dg-note {variable 'A\.[0-9]+' adjusted for OpenACC privatization level: 'gang'} "" { target { ! { openacc_host_selected || openacc_nvidia_accel_selected } } } l_compute$c_compute }
+    ! { dg-note {variable 'A\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: static} "" { target *-*-* } l_compute$c_compute }
     array = [(-2*i, i = 1, size(array))]
     !$acc loop gang private(array) ! { dg-line l_loop[incr c_loop] }
     ! { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop }
-- 
2.30.2