public inbox for gcc-patches@gcc.gnu.org
 help / color / mirror / Atom feed
From: Thomas Schwinge <thomas@codesourcery.com>
To: Julian Brown <julian@codesourcery.com>, <gcc-patches@gcc.gnu.org>
Cc: Andrew Stubbs <ams@codesourcery.com>, <fortran@gcc.gnu.org>
Subject: Add 'libgomp.oacc-c-c++-common/private-big-1.c' [PR105421] (was: amdgcn: Use FLAT addressing for all functions with pointer arguments [PR105421])
Date: Thu, 20 Oct 2022 12:19:28 +0200	[thread overview]
Message-ID: <87h6zyhk5r.fsf@euler.schwinge.homeip.net> (raw)
In-Reply-To: <87lepahkt3.fsf@euler.schwinge.homeip.net>

[-- 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


      reply	other threads:[~2022-10-20 10:19 UTC|newest]

Thread overview: 9+ messages / expand[flat|nested]  mbox.gz  Atom feed  top
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   ` Thomas Schwinge [this message]

Reply instructions:

You may reply publicly to this message via plain-text email
using any one of the following methods:

* Save the following mbox file, import it into your mail client,
  and reply-to-all from there: mbox

  Avoid top-posting and favor interleaved quoting:
  https://en.wikipedia.org/wiki/Posting_style#Interleaved_style

* Reply using the --to, --cc, and --in-reply-to
  switches of git-send-email(1):

  git send-email \
    --in-reply-to=87h6zyhk5r.fsf@euler.schwinge.homeip.net \
    --to=thomas@codesourcery.com \
    --cc=ams@codesourcery.com \
    --cc=fortran@gcc.gnu.org \
    --cc=gcc-patches@gcc.gnu.org \
    --cc=julian@codesourcery.com \
    /path/to/YOUR_REPLY

  https://kernel.org/pub/software/scm/git/docs/git-send-email.html

* If your mail client supports setting the In-Reply-To header
  via mailto: links, try the mailto: link
Be sure your reply has a Subject: header at the top and a blank line before the message body.
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).