From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: by sourceware.org (Postfix, from userid 1643) id A4C24385381C; Fri, 21 Oct 2022 12:59:33 +0000 (GMT) DKIM-Filter: OpenDKIM Filter v2.11.0 sourceware.org A4C24385381C DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=gcc.gnu.org; s=default; t=1666357173; bh=AKmcDs9oY/ozg6V3z+4YfItnehaJN4S9MfzpQltox9w=; h=From:To:Subject:Date:From; b=fXq3aKQuUtwxjrLbZyr3onPXCvVBGwEimhFqHZ1uGcgmTzUFkw2wMmEwUJtT1jLkG B4qIEF03YqC6ZN1+Yglerme0L63/jH7QZdNXuq3bSva21gHgj7U6pk9nhLZiqffTbm Pubk/ExeAmVSXw4vn5CiwR8oix6CTAy1RnS/X2qI= MIME-Version: 1.0 Content-Transfer-Encoding: 7bit Content-Type: text/plain; charset="utf-8" From: Thomas Schwinge To: gcc-cvs@gcc.gnu.org Subject: [gcc r12-8857] Add 'libgomp.oacc-c-c++-common/private-big-1.c' [PR105421] X-Act-Checkin: gcc X-Git-Author: Thomas Schwinge X-Git-Refname: refs/heads/releases/gcc-12 X-Git-Oldrev: dd2fdf114b59db1b0bfb0de497fc3ab11f034d89 X-Git-Newrev: b582496f205c1c3517d8415047fd78321f2c1cdf Message-Id: <20221021125933.A4C24385381C@sourceware.org> Date: Fri, 21 Oct 2022 12:59:33 +0000 (GMT) List-Id: https://gcc.gnu.org/g:b582496f205c1c3517d8415047fd78321f2c1cdf commit r12-8857-gb582496f205c1c3517d8415047fd78321f2c1cdf Author: Thomas Schwinge 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 +#include + + +/* 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; +}