public inbox for gcc-cvs@sourceware.org
help / color / mirror / Atom feed
* [gcc/devel/omp/gcc-12] Default compute dimensions (compile time)
@ 2022-06-29 14:34 Kwok Yeung
0 siblings, 0 replies; only message in thread
From: Kwok Yeung @ 2022-06-29 14:34 UTC (permalink / raw)
To: gcc-cvs
https://gcc.gnu.org/g:d9075a4c4782a7d95acc6ebee9aad6ce1eabbccb
commit d9075a4c4782a7d95acc6ebee9aad6ce1eabbccb
Author: Julian Brown <julian@codesourcery.com>
Date: Tue Feb 26 14:12:06 2019 -0800
Default compute dimensions (compile time)
Typo fix relative to last posted version.
2018-10-05 Nathan Sidwell <nathan@acm.org>
Tom de Vries <tdevries@suse.de>
Thomas Schwinge <thomas@codesourcery.com>
Julian Brown <julian@codesourcery.com>
gcc/
* doc/invoke.texi (fopenacc-dim): Update.
* omp-offload.cc (oacc_parse_default_dims): Update.
gcc/testsuite/
* c-c++-common/goacc/acc-icf.c: Update.
* c-c++-common/goacc/parallel-dims-1.c: Likewise.
* gfortran.dg/goacc/routine-4.f90: Likewise.
* gfortran.dg/goacc/routine-multiple-directives-1.f90: Likewise.
libgomp/
* testsuite/libgomp.oacc-c-c++-common/loop-default-compile.c: New.
* testsuite/libgomp.oacc-c-c++-common/loop-warn-1.c: New.
* testsuite/libgomp.oacc-c-c++-common/parallel-dims.c: Update.
Diff:
---
gcc/ChangeLog.omp | 8 +++++
gcc/doc/invoke.texi | 8 +++--
gcc/omp-offload.cc | 25 +++++++++------
gcc/testsuite/ChangeLog.omp | 10 ++++++
gcc/testsuite/c-c++-common/goacc/acc-icf.c | 4 +--
gcc/testsuite/c-c++-common/goacc/parallel-dims-1.c | 3 +-
gcc/testsuite/gfortran.dg/goacc/routine-4.f90 | 4 +++
.../goacc/routine-multiple-directives-1.f90 | 4 +--
libgomp/ChangeLog.omp | 9 ++++++
.../loop-default-compile.c | 13 ++++++++
.../libgomp.oacc-c-c++-common/loop-warn-1.c | 37 ++++++++++++++++++++++
11 files changed, 109 insertions(+), 16 deletions(-)
diff --git a/gcc/ChangeLog.omp b/gcc/ChangeLog.omp
index ec14b9c60a0..ea0d696347f 100644
--- a/gcc/ChangeLog.omp
+++ b/gcc/ChangeLog.omp
@@ -1,3 +1,11 @@
+2018-10-05 Nathan Sidwell <nathan@acm.org>
+ Tom de Vries <tdevries@suse.de>
+ Thomas Schwinge <thomas@codesourcery.com>
+ Julian Brown <julian@codesourcery.com>
+
+ * doc/invoke.texi (fopenacc-dim): Update.
+ * omp-offload.cc (oacc_parse_default_dims): Update.
+
2019-09-20 Chung-Lin Tang <cltang@codesourcery.com>
Cesar Philippidis <cesar@codesourcery.com>
diff --git a/gcc/doc/invoke.texi b/gcc/doc/invoke.texi
index 07b440190c3..3b19eb80212 100644
--- a/gcc/doc/invoke.texi
+++ b/gcc/doc/invoke.texi
@@ -2720,8 +2720,12 @@ have support for @option{-pthread}.
@cindex OpenACC accelerator programming
Specify default compute dimensions for parallel offload regions that do
not explicitly specify. The @var{geom} value is a triple of
-':'-separated sizes, in order 'gang', 'worker' and, 'vector'. A size
-can be omitted, to use a target-specific default value.
+':'-separated sizes, in order 'gang', 'worker' and, 'vector'. If a size
+is to be deferred until execution '-' can be used, alternatively a size
+can be omitted to use a target-specific default value. When deferring
+to runtime, the environment variable @var{GOMP_OPENACC_DIM} can be set.
+It has the same format as the option value, except that '-' is not
+permitted.
@item -fopenmp
@opindex fopenmp
diff --git a/gcc/omp-offload.cc b/gcc/omp-offload.cc
index 78c2982da5e..6131479b16f 100644
--- a/gcc/omp-offload.cc
+++ b/gcc/omp-offload.cc
@@ -845,8 +845,9 @@ oacc_get_min_dim (int dim)
}
/* Parse the default dimension parameter. This is a set of
- :-separated optional compute dimensions. Each specified dimension
- is a positive integer. When device type support is added, it is
+ :-separated optional compute dimensions. Each dimension is either
+ a positive integer, or '-' for a dynamic value computed at
+ runtime. When device type support is added, it is
planned to be a comma separated list of such compute dimensions,
with all but the first prefixed by the colon-terminated device
type. */
@@ -881,14 +882,20 @@ oacc_parse_default_dims (const char *dims)
if (*pos != ':')
{
- long val;
- const char *eptr;
+ long val = 0;
- errno = 0;
- val = strtol (pos, CONST_CAST (char **, &eptr), 10);
- if (errno || val <= 0 || (int) val != val)
- goto malformed;
- pos = eptr;
+ if (*pos == '-')
+ pos++;
+ else
+ {
+ const char *eptr;
+
+ errno = 0;
+ val = strtol (pos, CONST_CAST (char **, &eptr), 10);
+ if (errno || val <= 0 || (int) val != val)
+ goto malformed;
+ pos = eptr;
+ }
oacc_default_dims[ix] = (int) val;
}
}
diff --git a/gcc/testsuite/ChangeLog.omp b/gcc/testsuite/ChangeLog.omp
index 52910f0e422..4108c8dedc6 100644
--- a/gcc/testsuite/ChangeLog.omp
+++ b/gcc/testsuite/ChangeLog.omp
@@ -1,3 +1,13 @@
+2018-10-05 Nathan Sidwell <nathan@acm.org>
+ Tom de Vries <tdevries@suse.de>
+ Thomas Schwinge <thomas@codesourcery.com>
+ Julian Brown <julian@codesourcery.com>
+
+ * c-c++-common/goacc/acc-icf.c: Update.
+ * c-c++-common/goacc/parallel-dims-1.c: Likewise.
+ * gfortran.dg/goacc/routine-4.f90: Likewise.
+ * gfortran.dg/goacc/routine-multiple-directives-1.f90: Likewise.
+
2019-09-20 Chung-Lin Tang <cltang@codesourcery.com>
Cesar Philippidis <cesar@codesourcery.com>
diff --git a/gcc/testsuite/c-c++-common/goacc/acc-icf.c b/gcc/testsuite/c-c++-common/goacc/acc-icf.c
index 9cf119bf89c..bc2e0fd19b9 100644
--- a/gcc/testsuite/c-c++-common/goacc/acc-icf.c
+++ b/gcc/testsuite/c-c++-common/goacc/acc-icf.c
@@ -9,7 +9,7 @@
/* { dg-bogus "warning: region is worker partitioned but does not contain worker partitioned code" "TODO default 'gang' 'vector'" { xfail *-*-* } .+3 }
TODO It's the compiler's own decision to not use 'worker' parallelism here, so it doesn't make sense to bother the user about it. */
int
-routine1 (int n)
+routine1 (int n) /* { dg-bogus "region is worker partitioned but does not contain worker partitioned code" "" { xfail *-*-* } } */
{
int i;
@@ -24,7 +24,7 @@ routine1 (int n)
/* { dg-bogus "warning: region is worker partitioned but does not contain worker partitioned code" "TODO default 'gang' 'vector'" { xfail *-*-* } .+3 }
TODO It's the compiler's own decision to not use 'worker' parallelism here, so it doesn't make sense to bother the user about it. */
int
-routine2 (int n)
+routine2 (int n) /* { dg-bogus "region is worker partitioned but does not contain worker partitioned code" "" { xfail *-*-* } } */
{
int i;
diff --git a/gcc/testsuite/c-c++-common/goacc/parallel-dims-1.c b/gcc/testsuite/c-c++-common/goacc/parallel-dims-1.c
index 2a8d35d493d..6b1e7b22451 100644
--- a/gcc/testsuite/c-c++-common/goacc/parallel-dims-1.c
+++ b/gcc/testsuite/c-c++-common/goacc/parallel-dims-1.c
@@ -6,7 +6,8 @@
void f(int i)
{
-#pragma acc kernels num_gangs(i) num_workers(i) vector_length(i)
+#pragma acc kernels \
+ num_gangs(i) num_workers(i) vector_length(i)
;
#pragma acc parallel num_gangs(i) num_workers(i) vector_length(i)
diff --git a/gcc/testsuite/gfortran.dg/goacc/routine-4.f90 b/gcc/testsuite/gfortran.dg/goacc/routine-4.f90
index 53b1fbe5039..85fd50fb334 100644
--- a/gcc/testsuite/gfortran.dg/goacc/routine-4.f90
+++ b/gcc/testsuite/gfortran.dg/goacc/routine-4.f90
@@ -129,6 +129,7 @@ contains
integer, intent (inout) :: a(N)
integer :: i
+ !$acc loop gang worker vector
do i = 1, N
a(i) = a(i) - a(i)
end do
@@ -141,6 +142,7 @@ contains
integer, intent (inout) :: a(N)
integer :: i
+ !$acc loop worker vector
do i = 1, N
a(i) = a(i) - a(i)
end do
@@ -152,6 +154,7 @@ contains
integer, intent (inout) :: a(N)
integer :: i
+ !$acc loop vector
do i = 1, N
a(i) = a(i) - a(i)
end do
@@ -162,6 +165,7 @@ contains
integer, intent (inout) :: a(N)
integer :: i
+ !$acc loop seq
do i = 1, N
a(i) = a(i) - a(i)
end do
diff --git a/gcc/testsuite/gfortran.dg/goacc/routine-multiple-directives-1.f90 b/gcc/testsuite/gfortran.dg/goacc/routine-multiple-directives-1.f90
index 42bcb0e8d63..4249b404962 100644
--- a/gcc/testsuite/gfortran.dg/goacc/routine-multiple-directives-1.f90
+++ b/gcc/testsuite/gfortran.dg/goacc/routine-multiple-directives-1.f90
@@ -38,7 +38,7 @@
! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 's_2_nh' has 'nohost' clause" 1 "oaccloops" { target { ! offloading_enabled } } } }
! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 's_2_nh_' has 'nohost' clause" 1 "oaccloops" { target offloading_enabled } } }
- SUBROUTINE v_1
+ SUBROUTINE v_1 ! { dg-warning "region is vector partitioned but does not contain vector partitioned code" }
!$ACC ROUTINE VECTOR
!$ACC ROUTINE VECTOR
!$ACC ROUTINE(v_1) VECTOR
@@ -58,7 +58,7 @@
! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'v_1_nh' has 'nohost' clause" 1 "oaccloops" { target { ! offloading_enabled } } } }
! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'v_1_nh_' has 'nohost' clause" 1 "oaccloops" { target offloading_enabled } } }
- SUBROUTINE v_2
+ SUBROUTINE v_2 ! { dg-warning "region is vector partitioned but does not contain vector partitioned code" }
!$ACC ROUTINE(v_2) VECTOR
!$ACC ROUTINE VECTOR
!$ACC ROUTINE(v_2) VECTOR
diff --git a/libgomp/ChangeLog.omp b/libgomp/ChangeLog.omp
index 75345158736..3179b778c69 100644
--- a/libgomp/ChangeLog.omp
+++ b/libgomp/ChangeLog.omp
@@ -1,3 +1,12 @@
+2018-10-05 Nathan Sidwell <nathan@acm.org>
+ Tom de Vries <tdevries@suse.de>
+ Thomas Schwinge <thomas@codesourcery.com>
+ Julian Brown <julian@codesourcery.com>
+
+ * testsuite/libgomp.oacc-c-c++-common/loop-default-compile.c: New.
+ * testsuite/libgomp.oacc-c-c++-common/loop-warn-1.c: New.
+ * testsuite/libgomp.oacc-c-c++-common/parallel-dims.c: Likewise.
+
2018-10-22 James Norris <jnorris@codesourcery.com>
Cesar Philippidis <cesar@codesourcery.com>
Tom de Vries <tom@codesourcery.com>
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-default-compile.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-default-compile.c
new file mode 100644
index 00000000000..6c479e4eb25
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-default-compile.c
@@ -0,0 +1,13 @@
+/* { dg-additional-options "-fopenacc-dim=16:16" } */
+/* This code uses nvptx inline assembly guarded with acc_on_device, which is
+ not optimized away at -O0, and then confuses the target assembler.
+ { dg-skip-if "" { *-*-* } { "-O0" } { "" } } */
+/* { dg-set-target-env-var "GOMP_OPENACC_DIM" "8:8" } */
+
+#include "loop-default.h"
+
+int main ()
+{
+ /* Environment should be ignored. */
+ return test_1 (16, 16, 32);
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-warn-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-warn-1.c
new file mode 100644
index 00000000000..20a022f2758
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-warn-1.c
@@ -0,0 +1,37 @@
+
+/* Check warnings about suboptimal partitioning choices. */
+
+int main ()
+{
+ int ary[10];
+
+#pragma acc parallel copy(ary) num_gangs (1) /* { dg-warning "is not gang partitioned" } */
+ {
+ #pragma acc loop gang
+ for (int i = 0; i < 10; i++)
+ ary[i] = i;
+ }
+
+#pragma acc parallel copy(ary) num_workers (1) /* { dg-warning "is not worker partitioned" } */
+ {
+ #pragma acc loop worker
+ for (int i = 0; i < 10; i++)
+ ary[i] = i;
+ }
+
+#pragma acc parallel copy(ary) num_gangs (8) /* { dg-warning "is gang partitioned" } */
+ {
+ #pragma acc loop worker
+ for (int i = 0; i < 10; i++)
+ ary[i] = i;
+ }
+
+#pragma acc parallel copy(ary) num_workers (8) /* { dg-warning "is worker partitioned" } */
+ {
+ #pragma acc loop gang
+ for (int i = 0; i < 10; i++)
+ ary[i] = i;
+ }
+
+ return 0;
+}
^ permalink raw reply [flat|nested] only message in thread
only message in thread, other threads:[~2022-06-29 14:34 UTC | newest]
Thread overview: (only message) (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2022-06-29 14:34 [gcc/devel/omp/gcc-12] Default compute dimensions (compile time) Kwok Yeung
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).