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(-)
@@ -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;
@@ -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 } */
@@ -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 } */
@@ -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 } */
@@ -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 } */
@@ -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]];
@@ -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