public inbox for gcc-patches@gcc.gnu.org
 help / color / mirror / Atom feed
* [PATCH] [og12] amdgcn: Use FLAT addressing for all functions with pointer arguments
@ 2022-10-14 13:38 Julian Brown
  2022-10-14 13:38 ` [PATCH] [og12] OpenACC: Don't gang-privatize artificial variables Julian Brown
  2022-10-20 10:05 ` amdgcn: Use FLAT addressing for all functions with pointer arguments [PR105421] (was: [PATCH] [og12] amdgcn: Use FLAT addressing for all functions with pointer arguments) Thomas Schwinge
  0 siblings, 2 replies; 9+ messages in thread
From: Julian Brown @ 2022-10-14 13:38 UTC (permalink / raw)
  To: gcc-patches, Stubbs, Andrew, fortran

The GCN backend uses a heuristic to determine whether to use FLAT or
GLOBAL addressing in a particular (offload) function: namely, if a
function takes a pointer-to-scalar parameter, it is assumed that the
pointer may refer to "flat scratch" space, and thus FLAT addressing must
be used instead of GLOBAL.

I came up with this heuristic initially whilst working on support for
moving OpenACC gang-private variables into local-data share (scratch)
memory. The assumption that only scalar variables would be transformed in
that way turned out to be wrong.  For example, prior to the next patch in
the series, Fortran compiler-generated temporary structures were treated
as gang private and moved to LDS space, typically overflowing the region
allocated for such variables.  That will no longer happen after that
patch is applied, but there may be other cases of structs moving to LDS
space now or in the future that this patch may be needed for.

Tested with offloading to AMD GCN. I will apply shortly (to og12).

2022-10-14  Julian Brown  <julian@codesourcery.com>

gcc/
	* config/gcn/gcn.cc (gcn_detect_incoming_pointer_arg): Any pointer
	argument forces FLAT addressing mode, not just
	pointer-to-non-aggregate.
---
 gcc/ChangeLog.omp     |  6 ++++++
 gcc/config/gcn/gcn.cc | 15 +++++++++------
 2 files changed, 15 insertions(+), 6 deletions(-)

diff --git a/gcc/ChangeLog.omp b/gcc/ChangeLog.omp
index d296eb137e8..ceed4da9799 100644
--- a/gcc/ChangeLog.omp
+++ b/gcc/ChangeLog.omp
@@ -1,3 +1,9 @@
+2022-10-14  Julian Brown  <julian@codesourcery.com>
+
+	* config/gcn/gcn.cc (gcn_detect_incoming_pointer_arg): Any pointer
+	argument forces FLAT addressing mode, not just
+	pointer-to-non-aggregate.
+
 2022-10-12  Andrew Stubbs  <ams@codesourcery.com>
 
 	* config/gcn/gcn.cc (gcn_expand_builtin_1): Change gcn_full_exec_reg
diff --git a/gcc/config/gcn/gcn.cc b/gcc/config/gcn/gcn.cc
index 1f8d8e19971..b01131c0dc2 100644
--- a/gcc/config/gcn/gcn.cc
+++ b/gcc/config/gcn/gcn.cc
@@ -2819,10 +2819,14 @@ gcn_arg_partial_bytes (cumulative_args_t cum_v, const function_arg_info &arg)
   return (NUM_PARM_REGS - cum_num) * regsize;
 }
 
-/* A normal function which takes a pointer argument (to a scalar) may be
-   passed a pointer to LDS space (via a high-bits-set aperture), and that only
-   works with FLAT addressing, not GLOBAL.  Force FLAT addressing if the
-   function has an incoming pointer-to-scalar parameter.  */
+/* A normal function which takes a pointer argument may be passed a pointer to
+   LDS space (via a high-bits-set aperture), and that only works with FLAT
+   addressing, not GLOBAL.  Force FLAT addressing if the function has an
+   incoming pointer parameter.  NOTE: This is a heuristic that works in the
+   offloading case, but in general, a function might read global pointer
+   variables, etc. that may refer to LDS space or other special memory areas
+   not supported by GLOBAL instructions, and then this argument check would not
+   suffice.  */
 
 static void
 gcn_detect_incoming_pointer_arg (tree fndecl)
@@ -2832,8 +2836,7 @@ gcn_detect_incoming_pointer_arg (tree fndecl)
   for (tree arg = TYPE_ARG_TYPES (TREE_TYPE (fndecl));
        arg;
        arg = TREE_CHAIN (arg))
-    if (POINTER_TYPE_P (TREE_VALUE (arg))
-	&& !AGGREGATE_TYPE_P (TREE_TYPE (TREE_VALUE (arg))))
+    if (POINTER_TYPE_P (TREE_VALUE (arg)))
       cfun->machine->use_flat_addressing = true;
 }
 
-- 
2.29.2


^ permalink raw reply	[flat|nested] 9+ messages in thread

* [PATCH] [og12] OpenACC: Don't gang-privatize artificial variables
  2022-10-14 13:38 [PATCH] [og12] amdgcn: Use FLAT addressing for all functions with pointer arguments Julian Brown
@ 2022-10-14 13:38 ` Julian Brown
  2022-10-18 14:46   ` Thomas Schwinge
  2022-10-20 10:05 ` amdgcn: Use FLAT addressing for all functions with pointer arguments [PR105421] (was: [PATCH] [og12] amdgcn: Use FLAT addressing for all functions with pointer arguments) Thomas Schwinge
  1 sibling, 1 reply; 9+ messages in thread
From: Julian Brown @ 2022-10-14 13:38 UTC (permalink / raw)
  To: gcc-patches, Stubbs, Andrew, fortran

This patch prevents compiler-generated artificial variables from being
treated as privatization candidates for OpenACC.

The rationale is that e.g. "gang-private" variables actually must be
shared by each worker and vector spawned within a particular gang, but
that sharing is not necessary for any compiler-generated variable (at
least at present, but no such need is anticipated either).  Variables on
the stack (and machine registers) are already private per-"thread"
(gang, worker and/or vector), and that's fine for artificial variables.

Several tests need their scan output patterns adjusted to compensate.

Tested with offloading to AMD GCN. I will apply shortly (to og12).

2022-10-14  Julian Brown  <julian@codesourcery.com>

gcc/
	* omp-low.cc (oacc_privatization_candidate_p): Artificial vars are not
	privatization candidates.

libgomp/
	* testsuite/libgomp.oacc-fortran/declare-1.f90: Adjust scan output.
	* testsuite/libgomp.oacc-fortran/host_data-5.F90: Likewise.
	* testsuite/libgomp.oacc-fortran/if-1.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/print-1.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/privatized-ref-2.f90: Likewise.
---
 gcc/ChangeLog.omp                             |  5 ++++
 gcc/omp-low.cc                                | 22 +++++++++++++++
 libgomp/ChangeLog.omp                         |  8 ++++++
 .../libgomp.oacc-fortran/declare-1.f90        | 12 +++++++-
 .../libgomp.oacc-fortran/host_data-5.F90      | 28 +++++++++++--------
 .../testsuite/libgomp.oacc-fortran/if-1.f90   | 12 ++++----
 .../libgomp.oacc-fortran/print-1.f90          | 13 +--------
 .../libgomp.oacc-fortran/privatized-ref-2.f90 | 12 ++------
 8 files changed, 72 insertions(+), 40 deletions(-)

diff --git a/gcc/ChangeLog.omp b/gcc/ChangeLog.omp
index ceed4da9799..c34d0ec7c77 100644
--- a/gcc/ChangeLog.omp
+++ b/gcc/ChangeLog.omp
@@ -1,3 +1,8 @@
+2022-10-14  Julian Brown  <julian@codesourcery.com>
+
+	* omp-low.cc (oacc_privatization_candidate_p): Artificial vars are not
+	privatization candidates.
+
 2022-10-14  Julian Brown  <julian@codesourcery.com>
 
 	* config/gcn/gcn.cc (gcn_detect_incoming_pointer_arg): Any pointer
diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc
index d726eea2480..f171181e2c4 100644
--- a/gcc/omp-low.cc
+++ b/gcc/omp-low.cc
@@ -11400,6 +11400,28 @@ oacc_privatization_candidate_p (const location_t loc, const tree c,
 	}
     }
 
+  /* If an artificial variable has been added to a bind, e.g.
+     a compiler-generated temporary structure used by the Fortran front-end, do
+     not consider it as a privatization candidate.  Note that variables on
+     the stack are private per-thread by default: making them "gang-private"
+     for OpenACC actually means to share a single instance of a variable
+     amongst all workers and threads spawned within each gang.
+     At present, no compiler-generated artificial variables require such
+     sharing semantics, so this is safe.  */
+
+  if (res && DECL_ARTIFICIAL (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", "artificial");
+	}
+    }
+
   if (res)
     {
       if (dump_enabled_p ())
diff --git a/libgomp/ChangeLog.omp b/libgomp/ChangeLog.omp
index 7353fff2554..cb3541be378 100644
--- a/libgomp/ChangeLog.omp
+++ b/libgomp/ChangeLog.omp
@@ -1,3 +1,11 @@
+2022-10-14  Julian Brown  <julian@codesourcery.com>
+
+	* testsuite/libgomp.oacc-fortran/declare-1.f90: Adjust scan output.
+	* testsuite/libgomp.oacc-fortran/host_data-5.F90: Likewise.
+	* testsuite/libgomp.oacc-fortran/if-1.f90: Likewise.
+	* testsuite/libgomp.oacc-fortran/print-1.f90: Likewise.
+	* testsuite/libgomp.oacc-fortran/privatized-ref-2.f90: Likewise.
+
 2022-10-05  Tobias Burnus  <tobias@codesourcery.com>
 
 	Backport from mainline:
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/declare-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/declare-1.f90
index 51776a1d260..959e8941d5b 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/declare-1.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/declare-1.f90
@@ -25,6 +25,9 @@ module vars
 end module vars
 
 subroutine subr5 (a, b, c, d)
+  ! { dg-note {variable 'a\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-1 }
+  ! { dg-note {variable 'c\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-2 }
+  ! { dg-note {variable 'd\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-3 }
   implicit none
   integer, parameter :: N = 8
   integer :: i
@@ -51,6 +54,8 @@ subroutine subr5 (a, b, c, d)
 end subroutine
 
 subroutine subr4 (a, b)
+  ! { dg-note {variable 'a\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-1 }
+  ! { dg-note {variable 'b\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-2 }
   implicit none
   integer, parameter :: N = 8
   integer :: i
@@ -71,6 +76,8 @@ subroutine subr4 (a, b)
 end subroutine
 
 subroutine subr3 (a, c)
+  ! { dg-note {variable 'a\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-1 }
+  ! { dg-note {variable 'c\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-2 }
   implicit none
   integer, parameter :: N = 8
   integer :: i
@@ -92,6 +99,8 @@ subroutine subr3 (a, c)
 end subroutine
 
 subroutine subr2 (a, b, c)
+  ! { dg-note {variable 'a\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-1 }
+  ! { dg-note {variable 'c\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-2 }
   implicit none
   integer, parameter :: N = 8
   integer :: i
@@ -115,6 +124,7 @@ subroutine subr2 (a, b, c)
 end subroutine
 
 subroutine subr1 (a)
+  ! { dg-note {variable 'a\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-1 }
   implicit none
   integer, parameter :: N = 8
   integer :: i
@@ -215,7 +225,7 @@ program main
   ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } .-1 }
   ! { dg-note {variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-2 }
   ! { dg-note {variable 'S\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-3 }
-  ! { dg-note {variable 'desc\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} "TODO" { target *-*-* } .-4 }
+  ! { dg-note {variable 'desc\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target *-*-* } .-4 }
   use vars
   use openacc
   implicit none
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/host_data-5.F90 b/libgomp/testsuite/libgomp.oacc-fortran/host_data-5.F90
index 93e9ee09818..9b717929f02 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/host_data-5.F90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/host_data-5.F90
@@ -43,7 +43,7 @@ subroutine foo (p2, parr, host_p, host_parr, cond)
   ! { dg-note {variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target { ! openacc_host_selected } } .-2 }
   ! { dg-note {variable 'p\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-3 }
   ! { dg-note {variable 'parr\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-4 }
-  ! { dg-note {variable 'parm\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} "TODO" { target { ! openacc_host_selected } } .-5 }
+  ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target { ! openacc_host_selected } } .-5 }
 #if !ACC_MEM_SHARED
     if (acc_is_present(p, c_sizeof(p))) stop 5
     if (acc_is_present(parr, 1)) stop 6
@@ -54,9 +54,10 @@ subroutine foo (p2, parr, host_p, host_parr, cond)
     ! { dg-note {variable 'host_p\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-3 }
     ! { dg-note {variable 'parr\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-4 }
     ! { dg-note {variable 'host_parr\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-5 }
-    ! { dg-note {variable 'D\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} "TODO" { target *-*-* } .-6 }
-    ! { dg-note {variable 'transfer\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} "TODO" { target *-*-* } .-7 }
+    ! { dg-note {variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target *-*-* } .-6 }
+    ! { dg-note {variable 'transfer\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target *-*-* } .-7 }
     ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-8 }
+    ! { dg-note {variable 'p\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-9 }
       ! not mapped yet, so it will be equal to the host pointer.
       if (transfer(c_loc(p), host_p) /= host_p) stop 7
       if (transfer(c_loc(parr), host_parr) /= host_parr) stop 8
@@ -74,9 +75,9 @@ subroutine foo (p2, parr, host_p, host_parr, cond)
     ! { dg-note {variable 'host_p\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-5 }
     ! { dg-note {variable 'host_parr\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-6 }
     ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } .-7 }
-    ! { dg-note {variable 'parm\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} "TODO" { target *-*-* } .-8 }
-    ! { dg-note {variable 'D\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} "TODO" { target *-*-* } .-9 }
-    ! { dg-note {variable 'transfer\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} "TODO" { target *-*-* } .-10 }
+    ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target *-*-* } .-8 }
+    ! { dg-note {variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target *-*-* } .-9 }
+    ! { dg-note {variable 'transfer\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target *-*-* } .-10 }
     ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-11 }
       if (.not. acc_is_present(p, c_sizeof(p))) stop 11
       if (.not. acc_is_present(parr, 1)) stop 12
@@ -90,9 +91,10 @@ subroutine foo (p2, parr, host_p, host_parr, cond)
       ! { dg-note {variable 'host_p\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-3 }
       ! { dg-note {variable 'parr\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-4 }
       ! { dg-note {variable 'host_parr\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-5 }
-      ! { dg-note {variable 'D\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} "TODO" { target *-*-* } .-6 }
-      ! { dg-note {variable 'transfer\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} "TODO" { target *-*-* } .-7 }
+      ! { dg-note {variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target *-*-* } .-6 }
+      ! { dg-note {variable 'transfer\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target *-*-* } .-7 }
       ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-8 }
+      ! { dg-note {variable 'p\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-9 }
 #if ACC_MEM_SHARED
         if (transfer(c_loc(p), host_p) /= host_p) stop 15
         if (transfer(c_loc(parr), host_parr) /= host_parr) stop 16
@@ -110,8 +112,9 @@ subroutine foo (p2, parr, host_p, host_parr, cond)
         ! { dg-note {variable 'parr\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-4 }
         ! { dg-note {variable 'host_parr\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-5 }
         ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-6 }
-        ! { dg-note {variable 'D\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} "TODO" { target *-*-* } .-7 }
-        ! { dg-note {variable 'transfer\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} "TODO" { target *-*-* } .-8 }
+        ! { dg-note {variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target *-*-* } .-7 }
+        ! { dg-note {variable 'transfer\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target *-*-* } .-8 }
+	! { dg-note {variable 'p\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-9 }
 #if ACC_MEM_SHARED
         if (transfer(c_loc(p), host_p) /= host_p) stop 19
         if (transfer(c_loc(parr), host_parr) /= host_parr) stop 20
@@ -129,8 +132,9 @@ subroutine foo (p2, parr, host_p, host_parr, cond)
         ! { dg-note {variable 'parr\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-4 }
         ! { dg-note {variable 'host_parr\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-5 }
         ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-6 }
-        ! { dg-note {variable 'D\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} "TODO" { target *-*-* } .-7 }
-        ! { dg-note {variable 'transfer\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} "TODO" { target *-*-* } .-8 }
+        ! { dg-note {variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target *-*-* } .-7 }
+        ! { dg-note {variable 'transfer\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target *-*-* } .-8 }
+	! { dg-note {variable 'p\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-9 }
 #if ACC_MEM_SHARED
         if (transfer(c_loc(p), host_p) /= host_p) stop 23
         if (transfer(c_loc(parr), host_parr) /= host_parr) stop 24
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/if-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/if-1.f90
index c6d67647d4a..e0cfd912d0f 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/if-1.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/if-1.f90
@@ -382,7 +382,7 @@ program main
   b(:) = 1.0
 
   !$acc data copyin (a(1:N)) copyout (b(1:N)) if (0 == 1)
-  ! { dg-note {variable 'parm\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} "" { target { ! openacc_host_selected } } .-1 }
+  ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target { ! openacc_host_selected } } .-1 }
 
 #if !ACC_MEM_SHARED
   if (acc_is_present (a) .eqv. .TRUE.) STOP 21
@@ -396,7 +396,7 @@ program main
 
   !$acc data copyin (a(1:N)) if (1 == 1)
   ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-1 }
-  ! { dg-note {variable 'parm\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} "" { target { ! openacc_host_selected } } .-2 }
+  ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target { ! openacc_host_selected } } .-2 }
 
 #if !ACC_MEM_SHARED
     if (acc_is_present (a) .eqv. .FALSE.) STOP 23
@@ -404,7 +404,7 @@ program main
 
     !$acc data copyout (b(1:N)) if (0 == 1)
     ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-1 }
-    ! { dg-note {variable 'parm\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} "" { target { ! openacc_host_selected } } .-2 }
+    ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target { ! openacc_host_selected } } .-2 }
 #if !ACC_MEM_SHARED
       if (acc_is_present (b) .eqv. .TRUE.) STOP 24
 #endif
@@ -877,7 +877,7 @@ program main
   b(:) = 1.0
 
   !$acc data copyin (a(1:N)) copyout (b(1:N)) if (0 == 1)
-  ! { dg-note {variable 'parm\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} "" { target { ! openacc_host_selected } } .-1 }
+  ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target { ! openacc_host_selected } } .-1 }
 
 #if !ACC_MEM_SHARED
   if (acc_is_present (a) .eqv. .TRUE.) STOP 56
@@ -891,7 +891,7 @@ program main
 
   !$acc data copyin (a(1:N)) if (1 == 1)
   ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-1 }
-  ! { dg-note {variable 'parm\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} "" { target { ! openacc_host_selected } } .-2 }
+  ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target { ! openacc_host_selected } } .-2 }
 
 #if !ACC_MEM_SHARED
     if (acc_is_present (a) .eqv. .FALSE.) STOP 58
@@ -899,7 +899,7 @@ program main
 
     !$acc data copyout (b(1:N)) if (0 == 1)
     ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-1 }
-    ! { dg-note {variable 'parm\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} "" { target { ! openacc_host_selected } } .-2 }
+    ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target { ! openacc_host_selected } } .-2 }
 #if !ACC_MEM_SHARED
       if (acc_is_present (b) .eqv. .TRUE.) STOP 59
 #endif
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/print-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/print-1.f90
index 42a8538e1fb..d2f89d915f8 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/print-1.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/print-1.f90
@@ -6,15 +6,6 @@
 ! Separate file 'print-1-nvptx.f90' for nvptx offloading.
 ! { dg-skip-if "separate file" { offload_target_nvptx } }
 
-! For GCN offloading compilation, when gang-privatizing 'dt_parm.N'
-! (see below), we run into an 'gang-private data-share memory exhausted'
-! error: the default '-mgang-private-size' is too small.  Per
-! 'gcc/fortran/trans-io.cc'/'libgfortran/io/io.h', that one is
-! 'struct st_parameter_dt', which indeed is rather big.  Instead of
-! working out its exact size (which may vary per GCC configuration),
-! raise '-mgang-private-size' to an arbitrary high value.
-! { dg-additional-options "-foffload-options=amdgcn-amdhsa=-mgang-private-size=13579" { target openacc_radeon_accel_selected } }
-
 ! { dg-additional-options "-fopt-info-note-omp" }
 ! { dg-additional-options "-foffload=-fopt-info-note-omp" }
 
@@ -36,9 +27,7 @@ program main
   integer :: var = 42
 
 !$acc parallel ! { dg-line l_compute[incr c_compute] }
-  ! { dg-note {variable 'dt_parm\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l_compute$c_compute }
-  !   { dg-note {variable 'dt_parm\.[0-9]+' ought to be adjusted for OpenACC privatization level: 'gang'} {} { target *-*-* } l_compute$c_compute }
-  !   { dg-note {variable 'dt_parm\.[0-9]+' adjusted for OpenACC privatization level: 'gang'} {} { target { ! openacc_host_selected } } l_compute$c_compute }
+  ! { dg-note {variable 'dt_parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} {} { target *-*-* } l_compute$c_compute }
   write (0, '("The answer is ", I2)') var
 !$acc end parallel
 
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/privatized-ref-2.f90 b/libgomp/testsuite/libgomp.oacc-fortran/privatized-ref-2.f90
index b31f4066152..936285e9f69 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/privatized-ref-2.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/privatized-ref-2.f90
@@ -94,9 +94,7 @@ contains
     !$acc parallel copy(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 }
-    ! { dg-note {variable 'array\.[0-9]+' in 'private' clause is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_loop$c_loop }
-    ! { dg-note {variable 'array\.[0-9]+' ought to be adjusted for OpenACC privatization level: 'gang'} "" { target *-*-* } l_loop$c_loop }
-    ! { dg-note {variable 'array\.[0-9]+' adjusted for OpenACC privatization level: 'gang'} "" { target { ! { openacc_host_selected || { openacc_nvidia_accel_selected && __OPTIMIZE__ } } } } l_loop$c_loop }
+    ! { dg-note {variable 'array\.[0-9]+' in 'private' clause isn't candidate for adjusting OpenACC privatization level: artificial} "" { target *-*-* } l_loop$c_loop }
     ! { dg-message {sorry, unimplemented: target cannot support alloca} PR65181 { target openacc_nvidia_accel_selected } l_loop$c_loop }
     do i = 1, 10
       array(i) = 9*i
@@ -122,9 +120,7 @@ contains
     ! { dg-note {variable 'str' in 'private' clause is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_loop$c_loop }
     ! { dg-note {variable 'str' ought to be adjusted for OpenACC privatization level: 'gang'} "" { target *-*-* } l_loop$c_loop }
     ! { dg-note {variable 'str' adjusted for OpenACC privatization level: 'gang'} "" { target { ! { openacc_host_selected || { openacc_nvidia_accel_selected && __OPTIMIZE__ } } } } l_loop$c_loop }
-    ! { dg-note {variable 'char\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_loop$c_loop }
-    ! { dg-note {variable 'char\.[0-9]+' ought to be adjusted for OpenACC privatization level: 'gang'} "" { target *-*-* } l_loop$c_loop }
-    ! { dg-note {variable 'char\.[0-9]+' adjusted for OpenACC privatization level: 'gang'} "" { target { ! { openacc_host_selected || { openacc_nvidia_accel_selected && __OPTIMIZE__ } } } } l_loop$c_loop }
+    ! { dg-note {variable 'char\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target *-*-* } l_loop$c_loop }
     ! { dg-message {sorry, unimplemented: target cannot support alloca} PR65181 { target openacc_nvidia_accel_selected } l_loop$c_loop }
     do i = 1, 10
       str(i:i) = achar(ichar('A') + i)
@@ -167,9 +163,7 @@ contains
     ! { dg-note {variable 'scalar' in 'private' clause is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_loop$c_loop }
     ! { dg-note {variable 'scalar' ought to be adjusted for OpenACC privatization level: 'gang'} "" { target *-*-* } l_loop$c_loop }
     ! { dg-note {variable 'scalar' adjusted for OpenACC privatization level: 'gang'} "" { target { ! { openacc_host_selected || { openacc_nvidia_accel_selected && __OPTIMIZE__ } } } } l_loop$c_loop }
-    ! { dg-note {variable 'char\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_loop$c_loop }
-    ! { dg-note {variable 'char\.[0-9]+' ought to be adjusted for OpenACC privatization level: 'gang'} "" { target *-*-* } l_loop$c_loop }
-    ! { dg-note {variable 'char\.[0-9]+' adjusted for OpenACC privatization level: 'gang'} "" { target { ! { openacc_host_selected || { openacc_nvidia_accel_selected && __OPTIMIZE__ } } } } l_loop$c_loop }
+    ! { dg-note {variable 'char\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target *-*-* } l_loop$c_loop }
     do i = 1, 15
       scalar(i:i) = achar(ichar('A') + i)
     end do
-- 
2.29.2


^ permalink raw reply	[flat|nested] 9+ messages in thread

* Re: [PATCH] [og12] OpenACC: Don't gang-privatize artificial variables
  2022-10-14 13:38 ` [PATCH] [og12] OpenACC: Don't gang-privatize artificial variables Julian Brown
@ 2022-10-18 14:46   ` Thomas Schwinge
  2022-10-18 14:59     ` Julian Brown
  2022-10-28  8:51     ` OpenACC: Don't gang-privatize artificial variables [PR90115] " Thomas Schwinge
  0 siblings, 2 replies; 9+ messages in thread
From: Thomas Schwinge @ 2022-10-18 14:46 UTC (permalink / raw)
  To: Julian Brown; +Cc: gcc-patches, Andrew Stubbs, fortran

Hi Julian!

On 2022-10-14T13:38:56+0000, Julian Brown <julian@codesourcery.com> wrote:
> This patch prevents compiler-generated artificial variables from being
> treated as privatization candidates for OpenACC.
>
> The rationale is that e.g. "gang-private" variables actually must be
> shared by each worker and vector spawned within a particular gang, but
> that sharing is not necessary for any compiler-generated variable (at
> least at present, but no such need is anticipated either).  Variables on
> the stack (and machine registers) are already private per-"thread"
> (gang, worker and/or vector), and that's fine for artificial variables.

OK, that seems fine rationale for this change in behavior.
No contradicting test case jumped onto me, either.

> Several tests need their scan output patterns adjusted to compensate.

ACK -- surprisingly few.  (Some minor fine-tuning necessary for GCC
master branch, as had to be expected; I'm working on that.)

> --- a/gcc/omp-low.cc
> +++ b/gcc/omp-low.cc
> @@ -11400,6 +11400,28 @@ oacc_privatization_candidate_p (const location_t loc, const tree c,
>       }
>      }
>
> +  /* If an artificial variable has been added to a bind, e.g.
> +     a compiler-generated temporary structure used by the Fortran front-end, do
> +     not consider it as a privatization candidate.  Note that variables on
> +     the stack are private per-thread by default: making them "gang-private"
> +     for OpenACC actually means to share a single instance of a variable
> +     amongst all workers and threads spawned within each gang.
> +     At present, no compiler-generated artificial variables require such
> +     sharing semantics, so this is safe.  */
> +
> +  if (res && DECL_ARTIFICIAL (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", "artificial");
> +     }
> +    }

In the source code comment, you say "added to a bind", and that's indeed
what I was expecting, too, and thus put in:

       if (res && DECL_ARTIFICIAL (decl))
         {
    +      gcc_checking_assert (block);
    +
           res = false;

..., but to my surprised, that did fire in one occasion:

> --- a/libgomp/testsuite/libgomp.oacc-fortran/privatized-ref-2.f90
> +++ b/libgomp/testsuite/libgomp.oacc-fortran/privatized-ref-2.f90
> @@ -94,9 +94,7 @@ contains
>      !$acc parallel copy(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 }
> -    ! { dg-note {variable 'array\.[0-9]+' in 'private' clause is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_loop$c_loop }
> -    ! { dg-note {variable 'array\.[0-9]+' ought to be adjusted for OpenACC privatization level: 'gang'} "" { target *-*-* } l_loop$c_loop }
> -    ! { dg-note {variable 'array\.[0-9]+' adjusted for OpenACC privatization level: 'gang'} "" { target { ! { openacc_host_selected || { openacc_nvidia_accel_selected && __OPTIMIZE__ } } } } l_loop$c_loop }
> +    ! { dg-note {variable 'array\.[0-9]+' in 'private' clause isn't candidate for adjusting OpenACC privatization level: artificial} "" { target *-*-* } l_loop$c_loop }
>      ! { dg-message {sorry, unimplemented: target cannot support alloca} PR65181 { target openacc_nvidia_accel_selected } l_loop$c_loop }
>      do i = 1, 10
>        array(i) = 9*i

... here.  Note "variable 'array\.[0-9]+' in 'private' clause";
everywhere else we have "declared in block".

As part of your verification, have you already looked into whether the
new behavior is correct here, or does this one need to continue to be
"adjusted for OpenACC privatization level: 'gang'"?  If the latter,
should we check 'if (res && block && DECL_ARTIFICIAL (decl))' instead of
'if (res && DECL_ARTIFICIAL (decl))', or is there some wrong setting of
'DECL_ARTIFICIAL' -- or are we maybe looking at an inappropriate 'decl'?
(Thinking of commit r12-7580-g7a5e036b61aa088e6b8564bc9383d37dfbb4801e
"[OpenACC privatization] Analyze 'lookup_decl'-translated DECL [PR90115, PR102330, PR104774]",
for example.)


Grüße
 Thomas


> @@ -122,9 +120,7 @@ contains
>      ! { dg-note {variable 'str' in 'private' clause is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_loop$c_loop }
>      ! { dg-note {variable 'str' ought to be adjusted for OpenACC privatization level: 'gang'} "" { target *-*-* } l_loop$c_loop }
>      ! { dg-note {variable 'str' adjusted for OpenACC privatization level: 'gang'} "" { target { ! { openacc_host_selected || { openacc_nvidia_accel_selected && __OPTIMIZE__ } } } } l_loop$c_loop }
> -    ! { dg-note {variable 'char\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_loop$c_loop }
> -    ! { dg-note {variable 'char\.[0-9]+' ought to be adjusted for OpenACC privatization level: 'gang'} "" { target *-*-* } l_loop$c_loop }
> -    ! { dg-note {variable 'char\.[0-9]+' adjusted for OpenACC privatization level: 'gang'} "" { target { ! { openacc_host_selected || { openacc_nvidia_accel_selected && __OPTIMIZE__ } } } } l_loop$c_loop }
> +    ! { dg-note {variable 'char\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target *-*-* } l_loop$c_loop }
>      ! { dg-message {sorry, unimplemented: target cannot support alloca} PR65181 { target openacc_nvidia_accel_selected } l_loop$c_loop }
>      do i = 1, 10
>        str(i:i) = achar(ichar('A') + i)
> @@ -167,9 +163,7 @@ contains
>      ! { dg-note {variable 'scalar' in 'private' clause is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_loop$c_loop }
>      ! { dg-note {variable 'scalar' ought to be adjusted for OpenACC privatization level: 'gang'} "" { target *-*-* } l_loop$c_loop }
>      ! { dg-note {variable 'scalar' adjusted for OpenACC privatization level: 'gang'} "" { target { ! { openacc_host_selected || { openacc_nvidia_accel_selected && __OPTIMIZE__ } } } } l_loop$c_loop }
> -    ! { dg-note {variable 'char\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_loop$c_loop }
> -    ! { dg-note {variable 'char\.[0-9]+' ought to be adjusted for OpenACC privatization level: 'gang'} "" { target *-*-* } l_loop$c_loop }
> -    ! { dg-note {variable 'char\.[0-9]+' adjusted for OpenACC privatization level: 'gang'} "" { target { ! { openacc_host_selected || { openacc_nvidia_accel_selected && __OPTIMIZE__ } } } } l_loop$c_loop }
> +    ! { dg-note {variable 'char\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target *-*-* } l_loop$c_loop }
>      do i = 1, 15
>        scalar(i:i) = achar(ichar('A') + i)
>      end do
-----------------
Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955

^ permalink raw reply	[flat|nested] 9+ messages in thread

* Re: [PATCH] [og12] OpenACC: Don't gang-privatize artificial variables
  2022-10-18 14:46   ` Thomas Schwinge
@ 2022-10-18 14:59     ` Julian Brown
  2022-10-28  8:11       ` [og12] OpenACC: Don't gang-privatize artificial variables: restrict to blocks (was: [PATCH] [og12] OpenACC: Don't gang-privatize artificial variables) Thomas Schwinge
  2022-10-28  8:51     ` OpenACC: Don't gang-privatize artificial variables [PR90115] " Thomas Schwinge
  1 sibling, 1 reply; 9+ messages in thread
From: Julian Brown @ 2022-10-18 14:59 UTC (permalink / raw)
  To: Thomas Schwinge; +Cc: Andrew Stubbs, gcc-patches, fortran

On Tue, 18 Oct 2022 16:46:07 +0200
Thomas Schwinge <thomas@codesourcery.com> wrote:

> Hi Julian!
> 
> On 2022-10-14T13:38:56+0000, Julian Brown <julian@codesourcery.com>
> wrote:
> ..., but to my surprised, that did fire in one occasion:
> 
> > --- a/libgomp/testsuite/libgomp.oacc-fortran/privatized-ref-2.f90
> > +++ b/libgomp/testsuite/libgomp.oacc-fortran/privatized-ref-2.f90
> > @@ -94,9 +94,7 @@ contains
> >      !$acc parallel copy(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 }
> > -    ! { dg-note {variable 'array\.[0-9]+' in 'private' clause is
> > candidate for adjusting OpenACC privatization level} "" { target
> > *-*-* } l_loop$c_loop }
> > -    ! { dg-note {variable 'array\.[0-9]+' ought to be adjusted for
> > OpenACC privatization level: 'gang'} "" { target *-*-* }
> > l_loop$c_loop }
> > -    ! { dg-note {variable 'array\.[0-9]+' adjusted for OpenACC
> > privatization level: 'gang'} "" { target { ! {
> > openacc_host_selected || { openacc_nvidia_accel_selected &&
> > __OPTIMIZE__ } } } } l_loop$c_loop }
> > +    ! { dg-note {variable 'array\.[0-9]+' in 'private' clause
> > isn't candidate for adjusting OpenACC privatization level:
> > artificial} "" { target *-*-* } l_loop$c_loop } ! { dg-message
> > {sorry, unimplemented: target cannot support alloca} PR65181 {
> > target openacc_nvidia_accel_selected } l_loop$c_loop } do i = 1, 10
> > array(i) = 9*i  
> 
> ... here.  Note "variable 'array\.[0-9]+' in 'private' clause";
> everywhere else we have "declared in block".
> 
> As part of your verification, have you already looked into whether the
> new behavior is correct here, or does this one need to continue to be
> "adjusted for OpenACC privatization level: 'gang'"?  If the latter,
> should we check 'if (res && block && DECL_ARTIFICIAL (decl))' instead
> of 'if (res && DECL_ARTIFICIAL (decl))', or is there some wrong
> setting of 'DECL_ARTIFICIAL' -- or are we maybe looking at an
> inappropriate 'decl'? (Thinking of commit
> r12-7580-g7a5e036b61aa088e6b8564bc9383d37dfbb4801e "[OpenACC
> privatization] Analyze 'lookup_decl'-translated DECL [PR90115,
> PR102330, PR104774]", for example.)

I haven't looked in detail, but it seems to me that the "artificial"
flag isn't appropriate for that decl, which is (derived from?) a
user-visible symbol. So, I'm not sure what's going on there (and yes
the commit you mention looks like it could be relevant, I think?).
There are probably subtleties I'm not aware of...

HTH,

Julian


^ permalink raw reply	[flat|nested] 9+ messages in thread

* amdgcn: Use FLAT addressing for all functions with pointer arguments [PR105421] (was: [PATCH] [og12] amdgcn: Use FLAT addressing for all functions with pointer arguments)
  2022-10-14 13:38 [PATCH] [og12] amdgcn: Use FLAT addressing for all functions with pointer arguments Julian Brown
  2022-10-14 13:38 ` [PATCH] [og12] OpenACC: Don't gang-privatize artificial variables Julian Brown
@ 2022-10-20 10:05 ` Thomas Schwinge
  2022-10-20 10:19   ` Add 'libgomp.oacc-c-c++-common/private-big-1.c' [PR105421] (was: amdgcn: Use FLAT addressing for all functions with pointer arguments [PR105421]) Thomas Schwinge
  1 sibling, 1 reply; 9+ messages in thread
From: Thomas Schwinge @ 2022-10-20 10:05 UTC (permalink / raw)
  To: Julian Brown, gcc-patches; +Cc: Andrew Stubbs, fortran

[-- Attachment #1: Type: text/plain, Size: 1832 bytes --]

Hi!

On 2022-10-14T13:38:55+0000, Julian Brown <julian@codesourcery.com> wrote:
> The GCN backend uses a heuristic to determine whether to use FLAT or
> GLOBAL addressing in a particular (offload) function: namely, if a
> function takes a pointer-to-scalar parameter, it is assumed that the
> pointer may refer to "flat scratch" space, and thus FLAT addressing must
> be used instead of GLOBAL.
>
> I came up with this heuristic initially whilst working on support for
> moving OpenACC gang-private variables into local-data share (scratch)
> memory. The assumption that only scalar variables would be transformed in
> that way turned out to be wrong.  For example, prior to the next patch in
> the series, Fortran compiler-generated temporary structures were treated
> as gang private and moved to LDS space, typically overflowing the region
> allocated for such variables.  That will no longer happen after that
> patch is applied, but there may be other cases of structs moving to LDS
> space now or in the future that this patch may be needed for.
>
> Tested with offloading to AMD GCN. I will apply shortly (to og12).

Thanks.  I've verified that this does resolve PR105421
"GCN offloading, raised '-mgang-private-size': 'HSA_STATUS_ERROR_MEMORY_APERTURE_VIOLATION'"
and have thus added PR105421 tags to your commit log, and with that
pushed to master branch commit 7c55755d4c760de326809636531478fd7419e1e5
"amdgcn: Use FLAT addressing for all functions with pointer arguments [PR105421]",
see attached.


Grüße
 Thomas


-----------------
Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955

[-- Warning: decoded text below may be mangled, UTF-8 assumed --]
[-- Attachment #2: 0001-amdgcn-Use-FLAT-addressing-for-all-functions-with-po.patch --]
[-- Type: text/x-diff, Size: 3108 bytes --]

From 7c55755d4c760de326809636531478fd7419e1e5 Mon Sep 17 00:00:00 2001
From: Julian Brown <julian@codesourcery.com>
Date: Fri, 14 Oct 2022 11:06:07 +0000
Subject: [PATCH] amdgcn: Use FLAT addressing for all functions with pointer
 arguments [PR105421]

The GCN backend uses a heuristic to determine whether to use FLAT or
GLOBAL addressing in a particular (offload) function: namely, if a
function takes a pointer-to-scalar parameter, it is assumed that the
pointer may refer to "flat scratch" space, and thus FLAT addressing must
be used instead of GLOBAL.

I came up with this heuristic initially whilst working on support for
moving OpenACC gang-private variables into local-data share (scratch)
memory. The assumption that only scalar variables would be transformed in
that way turned out to be wrong.  For example, prior to the next patch in
the series, Fortran compiler-generated temporary structures were treated
as gang private and moved to LDS space, typically overflowing the region
allocated for such variables.  That will no longer happen after that
patch is applied, but there may be other cases of structs moving to LDS
space now or in the future that this patch may be needed for.

2022-10-14  Julian Brown  <julian@codesourcery.com>

	PR target/105421
gcc/
	* config/gcn/gcn.cc (gcn_detect_incoming_pointer_arg): Any pointer
	argument forces FLAT addressing mode, not just
	pointer-to-non-aggregate.
---
 gcc/config/gcn/gcn.cc | 15 +++++++++------
 1 file changed, 9 insertions(+), 6 deletions(-)

diff --git a/gcc/config/gcn/gcn.cc b/gcc/config/gcn/gcn.cc
index 8777255a5c6..a9ef5c3dc02 100644
--- a/gcc/config/gcn/gcn.cc
+++ b/gcc/config/gcn/gcn.cc
@@ -2809,10 +2809,14 @@ gcn_arg_partial_bytes (cumulative_args_t cum_v, const function_arg_info &arg)
   return (NUM_PARM_REGS - cum_num) * regsize;
 }
 
-/* A normal function which takes a pointer argument (to a scalar) may be
-   passed a pointer to LDS space (via a high-bits-set aperture), and that only
-   works with FLAT addressing, not GLOBAL.  Force FLAT addressing if the
-   function has an incoming pointer-to-scalar parameter.  */
+/* A normal function which takes a pointer argument may be passed a pointer to
+   LDS space (via a high-bits-set aperture), and that only works with FLAT
+   addressing, not GLOBAL.  Force FLAT addressing if the function has an
+   incoming pointer parameter.  NOTE: This is a heuristic that works in the
+   offloading case, but in general, a function might read global pointer
+   variables, etc. that may refer to LDS space or other special memory areas
+   not supported by GLOBAL instructions, and then this argument check would not
+   suffice.  */
 
 static void
 gcn_detect_incoming_pointer_arg (tree fndecl)
@@ -2822,8 +2826,7 @@ gcn_detect_incoming_pointer_arg (tree fndecl)
   for (tree arg = TYPE_ARG_TYPES (TREE_TYPE (fndecl));
        arg;
        arg = TREE_CHAIN (arg))
-    if (POINTER_TYPE_P (TREE_VALUE (arg))
-	&& !AGGREGATE_TYPE_P (TREE_TYPE (TREE_VALUE (arg))))
+    if (POINTER_TYPE_P (TREE_VALUE (arg)))
       cfun->machine->use_flat_addressing = true;
 }
 
-- 
2.35.1


^ permalink raw reply	[flat|nested] 9+ messages in thread

* Add 'libgomp.oacc-c-c++-common/private-big-1.c' [PR105421] (was: amdgcn: Use FLAT addressing for all functions with pointer arguments [PR105421])
  2022-10-20 10:05 ` amdgcn: Use FLAT addressing for all functions with pointer arguments [PR105421] (was: [PATCH] [og12] amdgcn: Use FLAT addressing for all functions with pointer arguments) Thomas Schwinge
@ 2022-10-20 10:19   ` Thomas Schwinge
  0 siblings, 0 replies; 9+ messages in thread
From: Thomas Schwinge @ 2022-10-20 10:19 UTC (permalink / raw)
  To: Julian Brown, gcc-patches; +Cc: Andrew Stubbs, fortran

[-- Attachment #1: Type: text/plain, Size: 2136 bytes --]

Hi!

On 2022-10-20T12:05:28+0200, I wrote:
> On 2022-10-14T13:38:55+0000, Julian Brown <julian@codesourcery.com> wrote:
>> The GCN backend uses a heuristic to determine whether to use FLAT or
>> GLOBAL addressing in a particular (offload) function: namely, if a
>> function takes a pointer-to-scalar parameter, it is assumed that the
>> pointer may refer to "flat scratch" space, and thus FLAT addressing must
>> be used instead of GLOBAL.
>>
>> I came up with this heuristic initially whilst working on support for
>> moving OpenACC gang-private variables into local-data share (scratch)
>> memory. The assumption that only scalar variables would be transformed in
>> that way turned out to be wrong.  For example, [...]
>> Fortran compiler-generated temporary structures were treated
>> as gang private and moved to LDS space, typically overflowing the region
>> allocated for such variables.  [...]
>> there may be other cases of structs moving to LDS
>> space now or in the future that this patch may be needed for.

When I (back then) had looked into PR105421
"GCN offloading, raised '-mgang-private-size': 'HSA_STATUS_ERROR_MEMORY_APERTURE_VIOLATION'",
I had been experimenting with different test codes, that all didn't
exhibit this problem.  Now I understand that 'struct' (as implied by
PR105421's Fortran 'write', for example) was the crucial thing there
(that is, 'AGGREGATE_TYPE_P (TREE_TYPE (TREE_VALUE (arg)))' in context of
the previous code).  With...

> pushed to master branch commit 7c55755d4c760de326809636531478fd7419e1e5
> "amdgcn: Use FLAT addressing for all functions with pointer arguments [PR105421]"

... that addressed, I've now pushed to master branch
commit c7ebee2378426eeca425ca5406af213a926f154c
"Add 'libgomp.oacc-c-c++-common/private-big-1.c' [PR105421]", see
attached.


Grüße
 Thomas


-----------------
Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955

[-- Warning: decoded text below may be mangled, UTF-8 assumed --]
[-- Attachment #2: 0001-Add-libgomp.oacc-c-c-common-private-big-1.c-PR105421.patch --]
[-- Type: text/x-diff, Size: 5367 bytes --]

From c7ebee2378426eeca425ca5406af213a926f154c Mon Sep 17 00:00:00 2001
From: Thomas Schwinge <thomas@codesourcery.com>
Date: Tue, 18 Oct 2022 00:13:47 +0200
Subject: [PATCH] Add 'libgomp.oacc-c-c++-common/private-big-1.c' [PR105421]

After commit r13-3404-g7c55755d4c760de326809636531478fd7419e1e5
"amdgcn: Use FLAT addressing for all functions with pointer arguments [PR105421]",
"big" private data now works for GCN offloading, too.

	PR target/105421
	libgomp/
	* testsuite/libgomp.oacc-c-c++-common/private-big-1.c: New.
---
 .../libgomp.oacc-c-c++-common/private-big-1.c | 100 ++++++++++++++++++
 1 file changed, 100 insertions(+)
 create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/private-big-1.c

diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/private-big-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/private-big-1.c
new file mode 100644
index 00000000000..c0e8db0c894
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/private-big-1.c
@@ -0,0 +1,100 @@
+/* Test "big" private data.  */
+
+/* { dg-additional-options -fno-inline } for stable results regarding OpenACC 'routine'.  */
+
+/* { dg-additional-options -fopt-info-all-omp }
+   { dg-additional-options --param=openacc-privatization=noisy }
+   { dg-additional-options -foffload=-fopt-info-all-omp }
+   { dg-additional-options -foffload=--param=openacc-privatization=noisy }
+   for testing/documenting aspects of that functionality.  */
+
+/* { dg-additional-options -Wopenacc-parallelism } for testing/documenting
+   aspects of that functionality.  */
+
+/* For GCN offloading compilation, we (expectedly) run into a
+   'gang-private data-share memory exhausted' error: the default
+   '-mgang-private-size' is too small.  Raise it so that 'uint32_t x[344]' plus
+   some internal-use data fits in:
+   { dg-additional-options -foffload-options=amdgcn-amdhsa=-mgang-private-size=1555 { target openacc_radeon_accel_selected } } */
+
+/* It's only with Tcl 8.5 (released in 2007) that "the variable 'varName'
+   passed to 'incr' may be unset, and in that case, it will be set to [...]",
+   so to maintain compatibility with earlier Tcl releases, we manually
+   initialize counter variables:
+   { dg-line l_dummy[variable c_compute 0 c_loop 0] }
+   { dg-message dummy {} { target iN-VAl-Id } l_dummy } to avoid
+   "WARNING: dg-line var l_dummy defined, but not used".  */
+
+#include <assert.h>
+#include <stdint.h>
+
+
+/* Based on 'private-variables.c:loop_g_5'.  */
+
+/* To demonstrate PR105421 "GCN offloading, raised '-mgang-private-size':
+   'HSA_STATUS_ERROR_MEMORY_APERTURE_VIOLATION'", a 'struct' indirection, for
+   example, has been necessary in combination with a separate routine.  */
+
+struct data
+{
+  uint32_t *x;
+  uint32_t *arr;
+  uint32_t i;
+};
+
+#pragma acc routine worker
+static void
+loop_g_5_r(struct data *data)
+{
+  uint32_t *x = data->x;
+  uint32_t *arr = data->arr;
+  uint32_t i = data->i;
+
+#pragma acc loop /* { dg-line l_loop[incr c_loop] } */
+  /* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop$c_loop } */
+  /* { dg-optimized {assigned OpenACC worker vector loop parallelism} {} { target *-*-* } l_loop$c_loop } */
+  for (int j = 0; j < 320; j++)
+    arr[i * 320 + j] += x[(i * 320 + j) % 344];
+}
+
+void loop_g_5()
+{
+  uint32_t x[344], i, arr[320 * 320];
+
+  for (i = 0; i < 320 * 320; i++)
+    arr[i] = i;
+
+  #pragma acc parallel copy(arr)
+  {
+    #pragma acc loop gang private(x) /* { dg-line l_loop[incr c_loop] } */
+    /* { dg-note {variable 'x' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l_loop$c_loop }
+       { dg-note {variable 'x' ought to be adjusted for OpenACC privatization level: 'gang'} {} { target *-*-* } l_loop$c_loop }
+       { dg-note {variable 'x' adjusted for OpenACC privatization level: 'gang'} {} { target { ! openacc_host_selected } } l_loop$c_loop } */
+    /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop$c_loop } */
+    /* { dg-note {variable 'data' declared in block is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l_loop$c_loop }
+       { dg-note {variable 'data' ought to be adjusted for OpenACC privatization level: 'gang'} {} { target *-*-* } l_loop$c_loop }
+       { dg-note {variable 'data' adjusted for OpenACC privatization level: 'gang'} {} { target { ! openacc_host_selected } } l_loop$c_loop } */
+    /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop$c_loop } */
+    /* { dg-optimized {assigned OpenACC gang loop parallelism} {} { target *-*-* } l_loop$c_loop } */
+    for (i = 0; i < 320; i++)
+      {
+        for (int j = 0; j < 344; j++)
+	  x[j] = j * (2 + i);
+
+	struct data data = { x, arr, i };
+	loop_g_5_r(&data); /* { dg-line l_compute[incr c_compute] } */
+	/* { dg-optimized {assigned OpenACC worker vector loop parallelism} {} { target *-*-* } l_compute$c_compute } */
+      }
+  }
+
+  for (i = 0; i < 320 * 320; i++)
+    assert(arr[i] == i + (i % 344) * (2 + (i / 320)));
+}
+
+
+int main ()
+{
+  loop_g_5();
+
+  return 0;
+}
-- 
2.35.1


^ permalink raw reply	[flat|nested] 9+ messages in thread

* [og12] OpenACC: Don't gang-privatize artificial variables: restrict to blocks (was: [PATCH] [og12] OpenACC: Don't gang-privatize artificial variables)
  2022-10-18 14:59     ` Julian Brown
@ 2022-10-28  8:11       ` Thomas Schwinge
  2022-10-28  8:20         ` Thomas Schwinge
  0 siblings, 1 reply; 9+ messages in thread
From: Thomas Schwinge @ 2022-10-28  8:11 UTC (permalink / raw)
  To: Julian Brown, gcc-patches, fortran; +Cc: Andrew Stubbs

Hi!

On 2022-10-18T15:59:24+0100, Julian Brown <julian@codesourcery.com> wrote:
> On Tue, 18 Oct 2022 16:46:07 +0200 Thomas Schwinge <thomas@codesourcery.com> wrote:
>> On 2022-10-14T13:38:56+0000, Julian Brown <julian@codesourcery.com> wrote:
>> ..., but to my surprised, that did fire in one occasion:
>>
>> > --- a/libgomp/testsuite/libgomp.oacc-fortran/privatized-ref-2.f90
>> > +++ b/libgomp/testsuite/libgomp.oacc-fortran/privatized-ref-2.f90
>> > @@ -94,9 +94,7 @@ contains
>> >      !$acc parallel copy(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 }
>> > -    ! { dg-note {variable 'array\.[0-9]+' in 'private' clause is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_loop$c_loop }
>> > -    ! { dg-note {variable 'array\.[0-9]+' ought to be adjusted for OpenACC privatization level: 'gang'} "" { target *-*-* } l_loop$c_loop }
>> > -    ! { dg-note {variable 'array\.[0-9]+' adjusted for OpenACC privatization level: 'gang'} "" { target { ! { openacc_host_selected || { openacc_nvidia_accel_selected && __OPTIMIZE__ } } } } l_loop$c_loop }
>> > +    ! { dg-note {variable 'array\.[0-9]+' in 'private' clause isn't candidate for adjusting OpenACC privatization level: artificial} "" { target *-*-* } l_loop$c_loop }
>> >      ! { dg-message {sorry, unimplemented: target cannot support alloca} PR65181 { target openacc_nvidia_accel_selected } l_loop$c_loop }
>> >      do i = 1, 10
>> >        array(i) = 9*i
>>
>> ... here.  Note "variable 'array\.[0-9]+' in 'private' clause";
>> everywhere else we have "declared in block".
>>
>> As part of your verification, have you already looked into whether the
>> new behavior is correct here, or does this one need to continue to be
>> "adjusted for OpenACC privatization level: 'gang'"?  If the latter,
>> should we check 'if (res && block && DECL_ARTIFICIAL (decl))' instead
>> of 'if (res && DECL_ARTIFICIAL (decl))', or is there some wrong
>> setting of 'DECL_ARTIFICIAL' -- or are we maybe looking at an
>> inappropriate 'decl'? (Thinking of commit
>> r12-7580-g7a5e036b61aa088e6b8564bc9383d37dfbb4801e "[OpenACC
>> privatization] Analyze 'lookup_decl'-translated DECL [PR90115,
>> PR102330, PR104774]", for example.)
>
> I haven't looked in detail, but it seems to me that the "artificial"
> flag isn't appropriate for that decl, which is (derived from?) a
> user-visible symbol. So, I'm not sure what's going on there (and yes
> the commit you mention looks like it could be relevant, I think?).
> There are probably subtleties I'm not aware of...

Until we've got that worked out, let's simply restrict the
'DECL_ARTIFICIAL' handling to 'block's only; pushed to devel/omp/gcc-12
commit 9a50d282f03f7f1e1ad00de917143a2a8e0c0ee0
"[og12] OpenACC: Don't gang-privatize artificial variables: restrict to blocks",
see attached.


Grüße
 Thomas
-----------------
Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955

^ permalink raw reply	[flat|nested] 9+ messages in thread

* Re: [og12] OpenACC: Don't gang-privatize artificial variables: restrict to blocks (was: [PATCH] [og12] OpenACC: Don't gang-privatize artificial variables)
  2022-10-28  8:11       ` [og12] OpenACC: Don't gang-privatize artificial variables: restrict to blocks (was: [PATCH] [og12] OpenACC: Don't gang-privatize artificial variables) Thomas Schwinge
@ 2022-10-28  8:20         ` Thomas Schwinge
  0 siblings, 0 replies; 9+ messages in thread
From: Thomas Schwinge @ 2022-10-28  8:20 UTC (permalink / raw)
  To: Julian Brown, gcc-patches, fortran; +Cc: Andrew Stubbs

[-- Attachment #1: Type: text/plain, Size: 3457 bytes --]

Hi!

On 2022-10-28T10:11:04+0200, I wrote:
> On 2022-10-18T15:59:24+0100, Julian Brown <julian@codesourcery.com> wrote:
>> On Tue, 18 Oct 2022 16:46:07 +0200 Thomas Schwinge <thomas@codesourcery.com> wrote:
>>> On 2022-10-14T13:38:56+0000, Julian Brown <julian@codesourcery.com> wrote:
>>> ..., but to my surprised, that did fire in one occasion:
>>>
>>> > --- a/libgomp/testsuite/libgomp.oacc-fortran/privatized-ref-2.f90
>>> > +++ b/libgomp/testsuite/libgomp.oacc-fortran/privatized-ref-2.f90
>>> > @@ -94,9 +94,7 @@ contains
>>> >      !$acc parallel copy(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 }
>>> > -    ! { dg-note {variable 'array\.[0-9]+' in 'private' clause is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_loop$c_loop }
>>> > -    ! { dg-note {variable 'array\.[0-9]+' ought to be adjusted for OpenACC privatization level: 'gang'} "" { target *-*-* } l_loop$c_loop }
>>> > -    ! { dg-note {variable 'array\.[0-9]+' adjusted for OpenACC privatization level: 'gang'} "" { target { ! { openacc_host_selected || { openacc_nvidia_accel_selected && __OPTIMIZE__ } } } } l_loop$c_loop }
>>> > +    ! { dg-note {variable 'array\.[0-9]+' in 'private' clause isn't candidate for adjusting OpenACC privatization level: artificial} "" { target *-*-* } l_loop$c_loop }
>>> >      ! { dg-message {sorry, unimplemented: target cannot support alloca} PR65181 { target openacc_nvidia_accel_selected } l_loop$c_loop }
>>> >      do i = 1, 10
>>> >        array(i) = 9*i
>>>
>>> ... here.  Note "variable 'array\.[0-9]+' in 'private' clause";
>>> everywhere else we have "declared in block".
>>>
>>> As part of your verification, have you already looked into whether the
>>> new behavior is correct here, or does this one need to continue to be
>>> "adjusted for OpenACC privatization level: 'gang'"?  If the latter,
>>> should we check 'if (res && block && DECL_ARTIFICIAL (decl))' instead
>>> of 'if (res && DECL_ARTIFICIAL (decl))', or is there some wrong
>>> setting of 'DECL_ARTIFICIAL' -- or are we maybe looking at an
>>> inappropriate 'decl'? (Thinking of commit
>>> r12-7580-g7a5e036b61aa088e6b8564bc9383d37dfbb4801e "[OpenACC
>>> privatization] Analyze 'lookup_decl'-translated DECL [PR90115,
>>> PR102330, PR104774]", for example.)
>>
>> I haven't looked in detail, but it seems to me that the "artificial"
>> flag isn't appropriate for that decl, which is (derived from?) a
>> user-visible symbol. So, I'm not sure what's going on there (and yes
>> the commit you mention looks like it could be relevant, I think?).
>> There are probably subtleties I'm not aware of...
>
> Until we've got that worked out, let's simply restrict the
> 'DECL_ARTIFICIAL' handling to 'block's only; pushed to devel/omp/gcc-12
> commit 9a50d282f03f7f1e1ad00de917143a2a8e0c0ee0
> "[og12] OpenACC: Don't gang-privatize artificial variables: restrict to blocks"

..., see attached now really.

Grüße
 Thomas


-----------------
Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955

[-- Warning: decoded text below may be mangled, UTF-8 assumed --]
[-- Attachment #2: 0001-og12-OpenACC-Don-t-gang-privatize-artificial-va.og12.patch --]
[-- Type: text/x-diff, Size: 2783 bytes --]

From 9a50d282f03f7f1e1ad00de917143a2a8e0c0ee0 Mon Sep 17 00:00:00 2001
From: Thomas Schwinge <thomas@codesourcery.com>
Date: Tue, 18 Oct 2022 16:59:54 +0200
Subject: [PATCH] [og12] OpenACC: Don't gang-privatize artificial variables:
 restrict to blocks

Follow-up to og12 commit d4504346d2a1d6ffecb8b2d8e3e04ab8ea259785
"[og12] OpenACC: Don't gang-privatize artificial variables", to restore
the previous behavior, until we understand what it means for a
'DECL_ARTIFICIAL' to appear in a 'private' clause.

	gcc/
	* omp-low.cc (oacc_privatization_candidate_p) <DECL_ARTIFICIAL>:
	Restrict to 'block's.
	libgomp/
	* testsuite/libgomp.oacc-fortran/privatized-ref-2.f90: Adjust.
---
 gcc/omp-low.cc                                              | 2 +-
 libgomp/testsuite/libgomp.oacc-fortran/privatized-ref-2.f90 | 4 +++-
 2 files changed, 4 insertions(+), 2 deletions(-)

diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc
index 002f91d930a..66aa11cd32d 100644
--- a/gcc/omp-low.cc
+++ b/gcc/omp-low.cc
@@ -11409,7 +11409,7 @@ oacc_privatization_candidate_p (const location_t loc, const tree c,
      At present, no compiler-generated artificial variables require such
      sharing semantics, so this is safe.  */
 
-  if (res && DECL_ARTIFICIAL (decl))
+  if (res && block && DECL_ARTIFICIAL (decl))
     {
       res = false;
 
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/privatized-ref-2.f90 b/libgomp/testsuite/libgomp.oacc-fortran/privatized-ref-2.f90
index 936285e9f69..498ef70b63a 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/privatized-ref-2.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/privatized-ref-2.f90
@@ -94,7 +94,9 @@ contains
     !$acc parallel copy(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 }
-    ! { dg-note {variable 'array\.[0-9]+' in 'private' clause isn't candidate for adjusting OpenACC privatization level: artificial} "" { target *-*-* } l_loop$c_loop }
+    ! { dg-note {variable 'array\.[0-9]+' in 'private' clause is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_loop$c_loop }
+    ! { dg-note {variable 'array\.[0-9]+' ought to be adjusted for OpenACC privatization level: 'gang'} "" { target *-*-* } l_loop$c_loop }
+    ! { dg-note {variable 'array\.[0-9]+' adjusted for OpenACC privatization level: 'gang'} "" { target { ! { openacc_host_selected || { openacc_nvidia_accel_selected && __OPTIMIZE__ } } } } l_loop$c_loop }
     ! { dg-message {sorry, unimplemented: target cannot support alloca} PR65181 { target openacc_nvidia_accel_selected } l_loop$c_loop }
     do i = 1, 10
       array(i) = 9*i
-- 
2.35.1


^ permalink raw reply	[flat|nested] 9+ messages in thread

* OpenACC: Don't gang-privatize artificial variables [PR90115] (was: [PATCH] [og12] OpenACC: Don't gang-privatize artificial variables)
  2022-10-18 14:46   ` Thomas Schwinge
  2022-10-18 14:59     ` Julian Brown
@ 2022-10-28  8:51     ` Thomas Schwinge
  1 sibling, 0 replies; 9+ messages in thread
From: Thomas Schwinge @ 2022-10-28  8:51 UTC (permalink / raw)
  To: Julian Brown, gcc-patches, fortran; +Cc: Andrew Stubbs

[-- Attachment #1: Type: text/plain, Size: 5438 bytes --]

Hi!

On 2022-10-18T16:46:07+0200, Thomas Schwinge <thomas@codesourcery.com> wrote:
> On 2022-10-14T13:38:56+0000, Julian Brown <julian@codesourcery.com> wrote:
>> This patch prevents compiler-generated artificial variables from being
>> treated as privatization candidates for OpenACC.
>>
>> The rationale is that e.g. "gang-private" variables actually must be
>> shared by each worker and vector spawned within a particular gang, but
>> that sharing is not necessary for any compiler-generated variable (at
>> least at present, but no such need is anticipated either).  Variables on
>> the stack (and machine registers) are already private per-"thread"
>> (gang, worker and/or vector), and that's fine for artificial variables.
>
> OK, that seems fine rationale for this change in behavior.
> No contradicting test case jumped onto me, either.

>> Several tests need their scan output patterns adjusted to compensate.
>
> ACK -- surprisingly few.  (Some minor fine-tuning necessary for GCC
> master branch, as had to be expected; I'm working on that.)

With those changes...

>> --- a/gcc/omp-low.cc
>> +++ b/gcc/omp-low.cc
>> @@ -11400,6 +11400,28 @@ oacc_privatization_candidate_p (const location_t loc, const tree c,
>>      }
>>      }
>>
>> +  /* If an artificial variable has been added to a bind, e.g.
>> +     a compiler-generated temporary structure used by the Fortran front-end, do
>> +     not consider it as a privatization candidate.  Note that variables on
>> +     the stack are private per-thread by default: making them "gang-private"
>> +     for OpenACC actually means to share a single instance of a variable
>> +     amongst all workers and threads spawned within each gang.
>> +     At present, no compiler-generated artificial variables require such
>> +     sharing semantics, so this is safe.  */
>> +
>> +  if (res && DECL_ARTIFICIAL (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", "artificial");
>> +    }
>> +    }
>
> In the source code comment, you say "added to a bind", and that's indeed
> what I was expecting, too, and thus put in:
>
>        if (res && DECL_ARTIFICIAL (decl))
>          {
>     +      gcc_checking_assert (block);
>     +
>            res = false;
>
> ..., but to my surprised, that did fire in one occasion:
>
>> --- a/libgomp/testsuite/libgomp.oacc-fortran/privatized-ref-2.f90
>> +++ b/libgomp/testsuite/libgomp.oacc-fortran/privatized-ref-2.f90
>> @@ -94,9 +94,7 @@ contains
>>      !$acc parallel copy(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 }
>> -    ! { dg-note {variable 'array\.[0-9]+' in 'private' clause is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_loop$c_loop }
>> -    ! { dg-note {variable 'array\.[0-9]+' ought to be adjusted for OpenACC privatization level: 'gang'} "" { target *-*-* } l_loop$c_loop }
>> -    ! { dg-note {variable 'array\.[0-9]+' adjusted for OpenACC privatization level: 'gang'} "" { target { ! { openacc_host_selected || { openacc_nvidia_accel_selected && __OPTIMIZE__ } } } } l_loop$c_loop }
>> +    ! { dg-note {variable 'array\.[0-9]+' in 'private' clause isn't candidate for adjusting OpenACC privatization level: artificial} "" { target *-*-* } l_loop$c_loop }
>>      ! { dg-message {sorry, unimplemented: target cannot support alloca} PR65181 { target openacc_nvidia_accel_selected } l_loop$c_loop }
>>      do i = 1, 10
>>        array(i) = 9*i
>
> ... here.  Note "variable 'array\.[0-9]+' in 'private' clause";
> everywhere else we have "declared in block".
>
> As part of your verification, have you already looked into whether the
> new behavior is correct here, or does this one need to continue to be
> "adjusted for OpenACC privatization level: 'gang'"?  If the latter,
> should we check 'if (res && block && DECL_ARTIFICIAL (decl))' instead of
> 'if (res && DECL_ARTIFICIAL (decl))'

..., and that change merged in, I've then pushed to master branch
commit 11e811d8e2f63667f60f73731bb934273f5882b8
"OpenACC: Don't gang-privatize artificial variables [PR90115]", see
attached.

Cherry-picked pushed to releases/gcc-12 branch in
commit 9b116c51a451995f1bae8fdac0748fcf3f06aafe
"OpenACC: Don't gang-privatize artificial variables [PR90115]", see
attached.

I've then also done a merge from releases/gcc-12 branch into
devel/omp/gcc-12 branch, taking care of merge conflicts due to
"fine-tuning necessary for GCC master branch", which shouldn't propagate
into devel/omp/gcc-12 branch.  Pushed to devel/omp/gcc-12 branch
commit 33eae55cd9effd9e0bb0c3659cc5dfc100b6fd4e
"Merge commit '9b116c51a451995f1bae8fdac0748fcf3f06aafe'".


Grüße
 Thomas


-----------------
Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955

[-- Warning: decoded text below may be mangled, UTF-8 assumed --]
[-- Attachment #2: 0001-OpenACC-Don-t-gang-privatize-artificial-variables-PR.patch --]
[-- Type: text/x-diff, Size: 21677 bytes --]

From 11e811d8e2f63667f60f73731bb934273f5882b8 Mon Sep 17 00:00:00 2001
From: Julian Brown <julian@codesourcery.com>
Date: Wed, 12 Oct 2022 20:44:57 +0000
Subject: [PATCH] OpenACC: Don't gang-privatize artificial variables [PR90115]

This patch prevents compiler-generated artificial variables from being
treated as privatization candidates for OpenACC.

The rationale is that e.g. "gang-private" variables actually must be
shared by each worker and vector spawned within a particular gang, but
that sharing is not necessary for any compiler-generated variable (at
least at present, but no such need is anticipated either).  Variables on
the stack (and machine registers) are already private per-"thread"
(gang, worker and/or vector), and that's fine for artificial variables.

We're restricting this to blocks, as we still need to understand what it
means for a 'DECL_ARTIFICIAL' to appear in a 'private' clause.

Several tests need their scan output patterns adjusted to compensate.

2022-10-14  Julian Brown  <julian@codesourcery.com>

	PR middle-end/90115
gcc/
	* omp-low.cc (oacc_privatization_candidate_p): Artificial vars are not
	privatization candidates.

libgomp/
	* testsuite/libgomp.oacc-fortran/declare-1.f90: Adjust scan output.
	* testsuite/libgomp.oacc-fortran/host_data-5.F90: Likewise.
	* testsuite/libgomp.oacc-fortran/if-1.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/print-1.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/privatized-ref-2.f90: Likewise.

Co-authored-by: Thomas Schwinge <thomas@codesourcery.com>
---
 gcc/omp-low.cc                                | 22 +++++++++++++++++
 .../libgomp.oacc-fortran/declare-1.f90        |  2 +-
 .../libgomp.oacc-fortran/host_data-5.F90      | 24 +++++++++----------
 .../testsuite/libgomp.oacc-fortran/if-1.f90   | 12 +++++-----
 .../libgomp.oacc-fortran/print-1.f90          | 13 +---------
 .../libgomp.oacc-fortran/privatized-ref-2.f90 |  8 ++-----
 6 files changed, 44 insertions(+), 37 deletions(-)

diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc
index a8809739001..82a93d00f67 100644
--- a/gcc/omp-low.cc
+++ b/gcc/omp-low.cc
@@ -10710,6 +10710,28 @@ oacc_privatization_candidate_p (const location_t loc, const tree c,
 	}
     }
 
+  /* If an artificial variable has been added to a bind, e.g.
+     a compiler-generated temporary structure used by the Fortran front-end, do
+     not consider it as a privatization candidate.  Note that variables on
+     the stack are private per-thread by default: making them "gang-private"
+     for OpenACC actually means to share a single instance of a variable
+     amongst all workers and threads spawned within each gang.
+     At present, no compiler-generated artificial variables require such
+     sharing semantics, so this is safe.  */
+
+  if (res && block && DECL_ARTIFICIAL (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", "artificial");
+	}
+    }
+
   if (res)
     {
       if (dump_enabled_p ())
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/declare-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/declare-1.f90
index 51776a1d260..89bd4a2d094 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/declare-1.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/declare-1.f90
@@ -215,7 +215,7 @@ program main
   ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } .-1 }
   ! { dg-note {variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-2 }
   ! { dg-note {variable 'S\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-3 }
-  ! { dg-note {variable 'desc\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} "TODO" { target *-*-* } .-4 }
+  ! { dg-note {variable 'desc\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target *-*-* } .-4 }
   use vars
   use openacc
   implicit none
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/host_data-5.F90 b/libgomp/testsuite/libgomp.oacc-fortran/host_data-5.F90
index 93e9ee09818..c3453a579ae 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/host_data-5.F90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/host_data-5.F90
@@ -43,7 +43,7 @@ subroutine foo (p2, parr, host_p, host_parr, cond)
   ! { dg-note {variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target { ! openacc_host_selected } } .-2 }
   ! { dg-note {variable 'p\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-3 }
   ! { dg-note {variable 'parr\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-4 }
-  ! { dg-note {variable 'parm\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} "TODO" { target { ! openacc_host_selected } } .-5 }
+  ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target { ! openacc_host_selected } } .-5 }
 #if !ACC_MEM_SHARED
     if (acc_is_present(p, c_sizeof(p))) stop 5
     if (acc_is_present(parr, 1)) stop 6
@@ -54,8 +54,8 @@ subroutine foo (p2, parr, host_p, host_parr, cond)
     ! { dg-note {variable 'host_p\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-3 }
     ! { dg-note {variable 'parr\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-4 }
     ! { dg-note {variable 'host_parr\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-5 }
-    ! { dg-note {variable 'D\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} "TODO" { target *-*-* } .-6 }
-    ! { dg-note {variable 'transfer\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} "TODO" { target *-*-* } .-7 }
+    ! { dg-note {variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target *-*-* } .-6 }
+    ! { dg-note {variable 'transfer\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target *-*-* } .-7 }
     ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-8 }
       ! not mapped yet, so it will be equal to the host pointer.
       if (transfer(c_loc(p), host_p) /= host_p) stop 7
@@ -74,9 +74,9 @@ subroutine foo (p2, parr, host_p, host_parr, cond)
     ! { dg-note {variable 'host_p\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-5 }
     ! { dg-note {variable 'host_parr\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-6 }
     ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } .-7 }
-    ! { dg-note {variable 'parm\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} "TODO" { target *-*-* } .-8 }
-    ! { dg-note {variable 'D\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} "TODO" { target *-*-* } .-9 }
-    ! { dg-note {variable 'transfer\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} "TODO" { target *-*-* } .-10 }
+    ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target *-*-* } .-8 }
+    ! { dg-note {variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target *-*-* } .-9 }
+    ! { dg-note {variable 'transfer\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target *-*-* } .-10 }
     ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-11 }
       if (.not. acc_is_present(p, c_sizeof(p))) stop 11
       if (.not. acc_is_present(parr, 1)) stop 12
@@ -90,8 +90,8 @@ subroutine foo (p2, parr, host_p, host_parr, cond)
       ! { dg-note {variable 'host_p\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-3 }
       ! { dg-note {variable 'parr\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-4 }
       ! { dg-note {variable 'host_parr\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-5 }
-      ! { dg-note {variable 'D\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} "TODO" { target *-*-* } .-6 }
-      ! { dg-note {variable 'transfer\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} "TODO" { target *-*-* } .-7 }
+      ! { dg-note {variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target *-*-* } .-6 }
+      ! { dg-note {variable 'transfer\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target *-*-* } .-7 }
       ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-8 }
 #if ACC_MEM_SHARED
         if (transfer(c_loc(p), host_p) /= host_p) stop 15
@@ -110,8 +110,8 @@ subroutine foo (p2, parr, host_p, host_parr, cond)
         ! { dg-note {variable 'parr\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-4 }
         ! { dg-note {variable 'host_parr\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-5 }
         ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-6 }
-        ! { dg-note {variable 'D\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} "TODO" { target *-*-* } .-7 }
-        ! { dg-note {variable 'transfer\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} "TODO" { target *-*-* } .-8 }
+        ! { dg-note {variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target *-*-* } .-7 }
+        ! { dg-note {variable 'transfer\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target *-*-* } .-8 }
 #if ACC_MEM_SHARED
         if (transfer(c_loc(p), host_p) /= host_p) stop 19
         if (transfer(c_loc(parr), host_parr) /= host_parr) stop 20
@@ -129,8 +129,8 @@ subroutine foo (p2, parr, host_p, host_parr, cond)
         ! { dg-note {variable 'parr\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-4 }
         ! { dg-note {variable 'host_parr\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-5 }
         ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-6 }
-        ! { dg-note {variable 'D\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} "TODO" { target *-*-* } .-7 }
-        ! { dg-note {variable 'transfer\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} "TODO" { target *-*-* } .-8 }
+        ! { dg-note {variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target *-*-* } .-7 }
+        ! { dg-note {variable 'transfer\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target *-*-* } .-8 }
 #if ACC_MEM_SHARED
         if (transfer(c_loc(p), host_p) /= host_p) stop 23
         if (transfer(c_loc(parr), host_parr) /= host_parr) stop 24
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/if-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/if-1.f90
index c6d67647d4a..e0cfd912d0f 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/if-1.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/if-1.f90
@@ -382,7 +382,7 @@ program main
   b(:) = 1.0
 
   !$acc data copyin (a(1:N)) copyout (b(1:N)) if (0 == 1)
-  ! { dg-note {variable 'parm\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} "" { target { ! openacc_host_selected } } .-1 }
+  ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target { ! openacc_host_selected } } .-1 }
 
 #if !ACC_MEM_SHARED
   if (acc_is_present (a) .eqv. .TRUE.) STOP 21
@@ -396,7 +396,7 @@ program main
 
   !$acc data copyin (a(1:N)) if (1 == 1)
   ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-1 }
-  ! { dg-note {variable 'parm\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} "" { target { ! openacc_host_selected } } .-2 }
+  ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target { ! openacc_host_selected } } .-2 }
 
 #if !ACC_MEM_SHARED
     if (acc_is_present (a) .eqv. .FALSE.) STOP 23
@@ -404,7 +404,7 @@ program main
 
     !$acc data copyout (b(1:N)) if (0 == 1)
     ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-1 }
-    ! { dg-note {variable 'parm\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} "" { target { ! openacc_host_selected } } .-2 }
+    ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target { ! openacc_host_selected } } .-2 }
 #if !ACC_MEM_SHARED
       if (acc_is_present (b) .eqv. .TRUE.) STOP 24
 #endif
@@ -877,7 +877,7 @@ program main
   b(:) = 1.0
 
   !$acc data copyin (a(1:N)) copyout (b(1:N)) if (0 == 1)
-  ! { dg-note {variable 'parm\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} "" { target { ! openacc_host_selected } } .-1 }
+  ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target { ! openacc_host_selected } } .-1 }
 
 #if !ACC_MEM_SHARED
   if (acc_is_present (a) .eqv. .TRUE.) STOP 56
@@ -891,7 +891,7 @@ program main
 
   !$acc data copyin (a(1:N)) if (1 == 1)
   ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-1 }
-  ! { dg-note {variable 'parm\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} "" { target { ! openacc_host_selected } } .-2 }
+  ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target { ! openacc_host_selected } } .-2 }
 
 #if !ACC_MEM_SHARED
     if (acc_is_present (a) .eqv. .FALSE.) STOP 58
@@ -899,7 +899,7 @@ program main
 
     !$acc data copyout (b(1:N)) if (0 == 1)
     ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-1 }
-    ! { dg-note {variable 'parm\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} "" { target { ! openacc_host_selected } } .-2 }
+    ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target { ! openacc_host_selected } } .-2 }
 #if !ACC_MEM_SHARED
       if (acc_is_present (b) .eqv. .TRUE.) STOP 59
 #endif
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/print-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/print-1.f90
index 42a8538e1fb..d2f89d915f8 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/print-1.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/print-1.f90
@@ -6,15 +6,6 @@
 ! Separate file 'print-1-nvptx.f90' for nvptx offloading.
 ! { dg-skip-if "separate file" { offload_target_nvptx } }
 
-! For GCN offloading compilation, when gang-privatizing 'dt_parm.N'
-! (see below), we run into an 'gang-private data-share memory exhausted'
-! error: the default '-mgang-private-size' is too small.  Per
-! 'gcc/fortran/trans-io.cc'/'libgfortran/io/io.h', that one is
-! 'struct st_parameter_dt', which indeed is rather big.  Instead of
-! working out its exact size (which may vary per GCC configuration),
-! raise '-mgang-private-size' to an arbitrary high value.
-! { dg-additional-options "-foffload-options=amdgcn-amdhsa=-mgang-private-size=13579" { target openacc_radeon_accel_selected } }
-
 ! { dg-additional-options "-fopt-info-note-omp" }
 ! { dg-additional-options "-foffload=-fopt-info-note-omp" }
 
@@ -36,9 +27,7 @@ program main
   integer :: var = 42
 
 !$acc parallel ! { dg-line l_compute[incr c_compute] }
-  ! { dg-note {variable 'dt_parm\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l_compute$c_compute }
-  !   { dg-note {variable 'dt_parm\.[0-9]+' ought to be adjusted for OpenACC privatization level: 'gang'} {} { target *-*-* } l_compute$c_compute }
-  !   { dg-note {variable 'dt_parm\.[0-9]+' adjusted for OpenACC privatization level: 'gang'} {} { target { ! openacc_host_selected } } l_compute$c_compute }
+  ! { dg-note {variable 'dt_parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} {} { target *-*-* } l_compute$c_compute }
   write (0, '("The answer is ", I2)') var
 !$acc end parallel
 
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/privatized-ref-2.f90 b/libgomp/testsuite/libgomp.oacc-fortran/privatized-ref-2.f90
index b31f4066152..498ef70b63a 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/privatized-ref-2.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/privatized-ref-2.f90
@@ -122,9 +122,7 @@ contains
     ! { dg-note {variable 'str' in 'private' clause is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_loop$c_loop }
     ! { dg-note {variable 'str' ought to be adjusted for OpenACC privatization level: 'gang'} "" { target *-*-* } l_loop$c_loop }
     ! { dg-note {variable 'str' adjusted for OpenACC privatization level: 'gang'} "" { target { ! { openacc_host_selected || { openacc_nvidia_accel_selected && __OPTIMIZE__ } } } } l_loop$c_loop }
-    ! { dg-note {variable 'char\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_loop$c_loop }
-    ! { dg-note {variable 'char\.[0-9]+' ought to be adjusted for OpenACC privatization level: 'gang'} "" { target *-*-* } l_loop$c_loop }
-    ! { dg-note {variable 'char\.[0-9]+' adjusted for OpenACC privatization level: 'gang'} "" { target { ! { openacc_host_selected || { openacc_nvidia_accel_selected && __OPTIMIZE__ } } } } l_loop$c_loop }
+    ! { dg-note {variable 'char\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target *-*-* } l_loop$c_loop }
     ! { dg-message {sorry, unimplemented: target cannot support alloca} PR65181 { target openacc_nvidia_accel_selected } l_loop$c_loop }
     do i = 1, 10
       str(i:i) = achar(ichar('A') + i)
@@ -167,9 +165,7 @@ contains
     ! { dg-note {variable 'scalar' in 'private' clause is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_loop$c_loop }
     ! { dg-note {variable 'scalar' ought to be adjusted for OpenACC privatization level: 'gang'} "" { target *-*-* } l_loop$c_loop }
     ! { dg-note {variable 'scalar' adjusted for OpenACC privatization level: 'gang'} "" { target { ! { openacc_host_selected || { openacc_nvidia_accel_selected && __OPTIMIZE__ } } } } l_loop$c_loop }
-    ! { dg-note {variable 'char\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_loop$c_loop }
-    ! { dg-note {variable 'char\.[0-9]+' ought to be adjusted for OpenACC privatization level: 'gang'} "" { target *-*-* } l_loop$c_loop }
-    ! { dg-note {variable 'char\.[0-9]+' adjusted for OpenACC privatization level: 'gang'} "" { target { ! { openacc_host_selected || { openacc_nvidia_accel_selected && __OPTIMIZE__ } } } } l_loop$c_loop }
+    ! { dg-note {variable 'char\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target *-*-* } l_loop$c_loop }
     do i = 1, 15
       scalar(i:i) = achar(ichar('A') + i)
     end do
-- 
2.35.1


[-- Warning: decoded text below may be mangled, UTF-8 assumed --]
[-- Attachment #3: 0001-OpenACC-Don-t-gang-privatize-artificial-variable.g12.patch --]
[-- Type: text/x-diff, Size: 21746 bytes --]

From 9b116c51a451995f1bae8fdac0748fcf3f06aafe Mon Sep 17 00:00:00 2001
From: Julian Brown <julian@codesourcery.com>
Date: Wed, 12 Oct 2022 20:44:57 +0000
Subject: [PATCH] OpenACC: Don't gang-privatize artificial variables [PR90115]

This patch prevents compiler-generated artificial variables from being
treated as privatization candidates for OpenACC.

The rationale is that e.g. "gang-private" variables actually must be
shared by each worker and vector spawned within a particular gang, but
that sharing is not necessary for any compiler-generated variable (at
least at present, but no such need is anticipated either).  Variables on
the stack (and machine registers) are already private per-"thread"
(gang, worker and/or vector), and that's fine for artificial variables.

We're restricting this to blocks, as we still need to understand what it
means for a 'DECL_ARTIFICIAL' to appear in a 'private' clause.

Several tests need their scan output patterns adjusted to compensate.

2022-10-14  Julian Brown  <julian@codesourcery.com>

	PR middle-end/90115
gcc/
	* omp-low.cc (oacc_privatization_candidate_p): Artificial vars are not
	privatization candidates.

libgomp/
	* testsuite/libgomp.oacc-fortran/declare-1.f90: Adjust scan output.
	* testsuite/libgomp.oacc-fortran/host_data-5.F90: Likewise.
	* testsuite/libgomp.oacc-fortran/if-1.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/print-1.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/privatized-ref-2.f90: Likewise.

Co-authored-by: Thomas Schwinge <thomas@codesourcery.com>
(cherry picked from commit 11e811d8e2f63667f60f73731bb934273f5882b8)
---
 gcc/omp-low.cc                                | 22 +++++++++++++++++
 .../libgomp.oacc-fortran/declare-1.f90        |  2 +-
 .../libgomp.oacc-fortran/host_data-5.F90      | 24 +++++++++----------
 .../testsuite/libgomp.oacc-fortran/if-1.f90   | 12 +++++-----
 .../libgomp.oacc-fortran/print-1.f90          | 13 +---------
 .../libgomp.oacc-fortran/privatized-ref-2.f90 |  8 ++-----
 6 files changed, 44 insertions(+), 37 deletions(-)

diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc
index e9482f2363b..a59996e2f13 100644
--- a/gcc/omp-low.cc
+++ b/gcc/omp-low.cc
@@ -10661,6 +10661,28 @@ oacc_privatization_candidate_p (const location_t loc, const tree c,
 	}
     }
 
+  /* If an artificial variable has been added to a bind, e.g.
+     a compiler-generated temporary structure used by the Fortran front-end, do
+     not consider it as a privatization candidate.  Note that variables on
+     the stack are private per-thread by default: making them "gang-private"
+     for OpenACC actually means to share a single instance of a variable
+     amongst all workers and threads spawned within each gang.
+     At present, no compiler-generated artificial variables require such
+     sharing semantics, so this is safe.  */
+
+  if (res && block && DECL_ARTIFICIAL (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", "artificial");
+	}
+    }
+
   if (res)
     {
       if (dump_enabled_p ())
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/declare-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/declare-1.f90
index 51776a1d260..89bd4a2d094 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/declare-1.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/declare-1.f90
@@ -215,7 +215,7 @@ program main
   ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } .-1 }
   ! { dg-note {variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-2 }
   ! { dg-note {variable 'S\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-3 }
-  ! { dg-note {variable 'desc\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} "TODO" { target *-*-* } .-4 }
+  ! { dg-note {variable 'desc\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target *-*-* } .-4 }
   use vars
   use openacc
   implicit none
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/host_data-5.F90 b/libgomp/testsuite/libgomp.oacc-fortran/host_data-5.F90
index 93e9ee09818..c3453a579ae 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/host_data-5.F90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/host_data-5.F90
@@ -43,7 +43,7 @@ subroutine foo (p2, parr, host_p, host_parr, cond)
   ! { dg-note {variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target { ! openacc_host_selected } } .-2 }
   ! { dg-note {variable 'p\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-3 }
   ! { dg-note {variable 'parr\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-4 }
-  ! { dg-note {variable 'parm\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} "TODO" { target { ! openacc_host_selected } } .-5 }
+  ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target { ! openacc_host_selected } } .-5 }
 #if !ACC_MEM_SHARED
     if (acc_is_present(p, c_sizeof(p))) stop 5
     if (acc_is_present(parr, 1)) stop 6
@@ -54,8 +54,8 @@ subroutine foo (p2, parr, host_p, host_parr, cond)
     ! { dg-note {variable 'host_p\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-3 }
     ! { dg-note {variable 'parr\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-4 }
     ! { dg-note {variable 'host_parr\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-5 }
-    ! { dg-note {variable 'D\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} "TODO" { target *-*-* } .-6 }
-    ! { dg-note {variable 'transfer\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} "TODO" { target *-*-* } .-7 }
+    ! { dg-note {variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target *-*-* } .-6 }
+    ! { dg-note {variable 'transfer\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target *-*-* } .-7 }
     ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-8 }
       ! not mapped yet, so it will be equal to the host pointer.
       if (transfer(c_loc(p), host_p) /= host_p) stop 7
@@ -74,9 +74,9 @@ subroutine foo (p2, parr, host_p, host_parr, cond)
     ! { dg-note {variable 'host_p\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-5 }
     ! { dg-note {variable 'host_parr\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-6 }
     ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } .-7 }
-    ! { dg-note {variable 'parm\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} "TODO" { target *-*-* } .-8 }
-    ! { dg-note {variable 'D\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} "TODO" { target *-*-* } .-9 }
-    ! { dg-note {variable 'transfer\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} "TODO" { target *-*-* } .-10 }
+    ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target *-*-* } .-8 }
+    ! { dg-note {variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target *-*-* } .-9 }
+    ! { dg-note {variable 'transfer\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target *-*-* } .-10 }
     ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-11 }
       if (.not. acc_is_present(p, c_sizeof(p))) stop 11
       if (.not. acc_is_present(parr, 1)) stop 12
@@ -90,8 +90,8 @@ subroutine foo (p2, parr, host_p, host_parr, cond)
       ! { dg-note {variable 'host_p\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-3 }
       ! { dg-note {variable 'parr\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-4 }
       ! { dg-note {variable 'host_parr\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-5 }
-      ! { dg-note {variable 'D\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} "TODO" { target *-*-* } .-6 }
-      ! { dg-note {variable 'transfer\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} "TODO" { target *-*-* } .-7 }
+      ! { dg-note {variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target *-*-* } .-6 }
+      ! { dg-note {variable 'transfer\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target *-*-* } .-7 }
       ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-8 }
 #if ACC_MEM_SHARED
         if (transfer(c_loc(p), host_p) /= host_p) stop 15
@@ -110,8 +110,8 @@ subroutine foo (p2, parr, host_p, host_parr, cond)
         ! { dg-note {variable 'parr\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-4 }
         ! { dg-note {variable 'host_parr\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-5 }
         ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-6 }
-        ! { dg-note {variable 'D\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} "TODO" { target *-*-* } .-7 }
-        ! { dg-note {variable 'transfer\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} "TODO" { target *-*-* } .-8 }
+        ! { dg-note {variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target *-*-* } .-7 }
+        ! { dg-note {variable 'transfer\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target *-*-* } .-8 }
 #if ACC_MEM_SHARED
         if (transfer(c_loc(p), host_p) /= host_p) stop 19
         if (transfer(c_loc(parr), host_parr) /= host_parr) stop 20
@@ -129,8 +129,8 @@ subroutine foo (p2, parr, host_p, host_parr, cond)
         ! { dg-note {variable 'parr\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-4 }
         ! { dg-note {variable 'host_parr\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-5 }
         ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-6 }
-        ! { dg-note {variable 'D\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} "TODO" { target *-*-* } .-7 }
-        ! { dg-note {variable 'transfer\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} "TODO" { target *-*-* } .-8 }
+        ! { dg-note {variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target *-*-* } .-7 }
+        ! { dg-note {variable 'transfer\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target *-*-* } .-8 }
 #if ACC_MEM_SHARED
         if (transfer(c_loc(p), host_p) /= host_p) stop 23
         if (transfer(c_loc(parr), host_parr) /= host_parr) stop 24
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/if-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/if-1.f90
index c6d67647d4a..e0cfd912d0f 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/if-1.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/if-1.f90
@@ -382,7 +382,7 @@ program main
   b(:) = 1.0
 
   !$acc data copyin (a(1:N)) copyout (b(1:N)) if (0 == 1)
-  ! { dg-note {variable 'parm\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} "" { target { ! openacc_host_selected } } .-1 }
+  ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target { ! openacc_host_selected } } .-1 }
 
 #if !ACC_MEM_SHARED
   if (acc_is_present (a) .eqv. .TRUE.) STOP 21
@@ -396,7 +396,7 @@ program main
 
   !$acc data copyin (a(1:N)) if (1 == 1)
   ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-1 }
-  ! { dg-note {variable 'parm\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} "" { target { ! openacc_host_selected } } .-2 }
+  ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target { ! openacc_host_selected } } .-2 }
 
 #if !ACC_MEM_SHARED
     if (acc_is_present (a) .eqv. .FALSE.) STOP 23
@@ -404,7 +404,7 @@ program main
 
     !$acc data copyout (b(1:N)) if (0 == 1)
     ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-1 }
-    ! { dg-note {variable 'parm\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} "" { target { ! openacc_host_selected } } .-2 }
+    ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target { ! openacc_host_selected } } .-2 }
 #if !ACC_MEM_SHARED
       if (acc_is_present (b) .eqv. .TRUE.) STOP 24
 #endif
@@ -877,7 +877,7 @@ program main
   b(:) = 1.0
 
   !$acc data copyin (a(1:N)) copyout (b(1:N)) if (0 == 1)
-  ! { dg-note {variable 'parm\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} "" { target { ! openacc_host_selected } } .-1 }
+  ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target { ! openacc_host_selected } } .-1 }
 
 #if !ACC_MEM_SHARED
   if (acc_is_present (a) .eqv. .TRUE.) STOP 56
@@ -891,7 +891,7 @@ program main
 
   !$acc data copyin (a(1:N)) if (1 == 1)
   ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-1 }
-  ! { dg-note {variable 'parm\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} "" { target { ! openacc_host_selected } } .-2 }
+  ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target { ! openacc_host_selected } } .-2 }
 
 #if !ACC_MEM_SHARED
     if (acc_is_present (a) .eqv. .FALSE.) STOP 58
@@ -899,7 +899,7 @@ program main
 
     !$acc data copyout (b(1:N)) if (0 == 1)
     ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-1 }
-    ! { dg-note {variable 'parm\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} "" { target { ! openacc_host_selected } } .-2 }
+    ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target { ! openacc_host_selected } } .-2 }
 #if !ACC_MEM_SHARED
       if (acc_is_present (b) .eqv. .TRUE.) STOP 59
 #endif
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/print-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/print-1.f90
index 42a8538e1fb..d2f89d915f8 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/print-1.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/print-1.f90
@@ -6,15 +6,6 @@
 ! Separate file 'print-1-nvptx.f90' for nvptx offloading.
 ! { dg-skip-if "separate file" { offload_target_nvptx } }
 
-! For GCN offloading compilation, when gang-privatizing 'dt_parm.N'
-! (see below), we run into an 'gang-private data-share memory exhausted'
-! error: the default '-mgang-private-size' is too small.  Per
-! 'gcc/fortran/trans-io.cc'/'libgfortran/io/io.h', that one is
-! 'struct st_parameter_dt', which indeed is rather big.  Instead of
-! working out its exact size (which may vary per GCC configuration),
-! raise '-mgang-private-size' to an arbitrary high value.
-! { dg-additional-options "-foffload-options=amdgcn-amdhsa=-mgang-private-size=13579" { target openacc_radeon_accel_selected } }
-
 ! { dg-additional-options "-fopt-info-note-omp" }
 ! { dg-additional-options "-foffload=-fopt-info-note-omp" }
 
@@ -36,9 +27,7 @@ program main
   integer :: var = 42
 
 !$acc parallel ! { dg-line l_compute[incr c_compute] }
-  ! { dg-note {variable 'dt_parm\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l_compute$c_compute }
-  !   { dg-note {variable 'dt_parm\.[0-9]+' ought to be adjusted for OpenACC privatization level: 'gang'} {} { target *-*-* } l_compute$c_compute }
-  !   { dg-note {variable 'dt_parm\.[0-9]+' adjusted for OpenACC privatization level: 'gang'} {} { target { ! openacc_host_selected } } l_compute$c_compute }
+  ! { dg-note {variable 'dt_parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} {} { target *-*-* } l_compute$c_compute }
   write (0, '("The answer is ", I2)') var
 !$acc end parallel
 
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/privatized-ref-2.f90 b/libgomp/testsuite/libgomp.oacc-fortran/privatized-ref-2.f90
index b31f4066152..498ef70b63a 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/privatized-ref-2.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/privatized-ref-2.f90
@@ -122,9 +122,7 @@ contains
     ! { dg-note {variable 'str' in 'private' clause is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_loop$c_loop }
     ! { dg-note {variable 'str' ought to be adjusted for OpenACC privatization level: 'gang'} "" { target *-*-* } l_loop$c_loop }
     ! { dg-note {variable 'str' adjusted for OpenACC privatization level: 'gang'} "" { target { ! { openacc_host_selected || { openacc_nvidia_accel_selected && __OPTIMIZE__ } } } } l_loop$c_loop }
-    ! { dg-note {variable 'char\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_loop$c_loop }
-    ! { dg-note {variable 'char\.[0-9]+' ought to be adjusted for OpenACC privatization level: 'gang'} "" { target *-*-* } l_loop$c_loop }
-    ! { dg-note {variable 'char\.[0-9]+' adjusted for OpenACC privatization level: 'gang'} "" { target { ! { openacc_host_selected || { openacc_nvidia_accel_selected && __OPTIMIZE__ } } } } l_loop$c_loop }
+    ! { dg-note {variable 'char\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target *-*-* } l_loop$c_loop }
     ! { dg-message {sorry, unimplemented: target cannot support alloca} PR65181 { target openacc_nvidia_accel_selected } l_loop$c_loop }
     do i = 1, 10
       str(i:i) = achar(ichar('A') + i)
@@ -167,9 +165,7 @@ contains
     ! { dg-note {variable 'scalar' in 'private' clause is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_loop$c_loop }
     ! { dg-note {variable 'scalar' ought to be adjusted for OpenACC privatization level: 'gang'} "" { target *-*-* } l_loop$c_loop }
     ! { dg-note {variable 'scalar' adjusted for OpenACC privatization level: 'gang'} "" { target { ! { openacc_host_selected || { openacc_nvidia_accel_selected && __OPTIMIZE__ } } } } l_loop$c_loop }
-    ! { dg-note {variable 'char\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_loop$c_loop }
-    ! { dg-note {variable 'char\.[0-9]+' ought to be adjusted for OpenACC privatization level: 'gang'} "" { target *-*-* } l_loop$c_loop }
-    ! { dg-note {variable 'char\.[0-9]+' adjusted for OpenACC privatization level: 'gang'} "" { target { ! { openacc_host_selected || { openacc_nvidia_accel_selected && __OPTIMIZE__ } } } } l_loop$c_loop }
+    ! { dg-note {variable 'char\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target *-*-* } l_loop$c_loop }
     do i = 1, 15
       scalar(i:i) = achar(ichar('A') + i)
     end do
-- 
2.35.1


^ permalink raw reply	[flat|nested] 9+ messages in thread

end of thread, other threads:[~2022-10-28  8:51 UTC | newest]

Thread overview: 9+ messages (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2022-10-14 13:38 [PATCH] [og12] amdgcn: Use FLAT addressing for all functions with pointer arguments Julian Brown
2022-10-14 13:38 ` [PATCH] [og12] OpenACC: Don't gang-privatize artificial variables Julian Brown
2022-10-18 14:46   ` Thomas Schwinge
2022-10-18 14:59     ` Julian Brown
2022-10-28  8:11       ` [og12] OpenACC: Don't gang-privatize artificial variables: restrict to blocks (was: [PATCH] [og12] OpenACC: Don't gang-privatize artificial variables) Thomas Schwinge
2022-10-28  8:20         ` Thomas Schwinge
2022-10-28  8:51     ` OpenACC: Don't gang-privatize artificial variables [PR90115] " Thomas Schwinge
2022-10-20 10:05 ` amdgcn: Use FLAT addressing for all functions with pointer arguments [PR105421] (was: [PATCH] [og12] amdgcn: Use FLAT addressing for all functions with pointer arguments) Thomas Schwinge
2022-10-20 10:19   ` Add 'libgomp.oacc-c-c++-common/private-big-1.c' [PR105421] (was: amdgcn: Use FLAT addressing for all functions with pointer arguments [PR105421]) Thomas Schwinge

This is a public inbox, see mirroring instructions
for how to clone and mirror all data and code used for this inbox;
as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).