public inbox for gcc-cvs@sourceware.org
help / color / mirror / Atom feed
* [gcc r12-8857] Add 'libgomp.oacc-c-c++-common/private-big-1.c' [PR105421]
@ 2022-10-21 12:59 Thomas Schwinge
0 siblings, 0 replies; only message in thread
From: Thomas Schwinge @ 2022-10-21 12:59 UTC (permalink / raw)
To: gcc-cvs
https://gcc.gnu.org/g:b582496f205c1c3517d8415047fd78321f2c1cdf
commit r12-8857-gb582496f205c1c3517d8415047fd78321f2c1cdf
Author: Thomas Schwinge <thomas@codesourcery.com>
Date: Tue Oct 18 00:13:47 2022 +0200
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.
(cherry picked from commit c7ebee2378426eeca425ca5406af213a926f154c)
Diff:
---
.../libgomp.oacc-c-c++-common/private-big-1.c | 100 +++++++++++++++++++++
1 file changed, 100 insertions(+)
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;
+}
^ permalink raw reply [flat|nested] only message in thread
only message in thread, other threads:[~2022-10-21 12:59 UTC | newest]
Thread overview: (only message) (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2022-10-21 12:59 [gcc r12-8857] Add 'libgomp.oacc-c-c++-common/private-big-1.c' [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).