https://gcc.gnu.org/g:48aca3e3007df5918dc8edd29c42adf3701497da
commit 48aca3e3007df5918dc8edd29c42adf3701497da
Author: Julian Brown <jul...@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 <nat...@acm.org>
Tom de Vries <tdevr...@suse.de>
Thomas Schwinge <tho...@codesourcery.com>
Julian Brown <jul...@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 4ba97e30d5c..952a79c8c24 100644
--- a/gcc/ChangeLog.omp
+++ b/gcc/ChangeLog.omp
@@ -1,3 +1,11 @@
+2018-10-05 Nathan Sidwell <nat...@acm.org>
+ Tom de Vries <tdevr...@suse.de>
+ Thomas Schwinge <tho...@codesourcery.com>
+ Julian Brown <jul...@codesourcery.com>
+
+ * doc/invoke.texi (fopenacc-dim): Update.
+ * omp-offload.cc (oacc_parse_default_dims): Update.
+
2018-10-30 Cesar Philippidis <ce...@codesourcery.com>
* config/nvptx/nvptx.cc (nvptx_propagate_unified): New.
diff --git a/gcc/doc/invoke.texi b/gcc/doc/invoke.texi
index 9456ced468a..3e2f9bd5b4b 100644
--- a/gcc/doc/invoke.texi
+++ b/gcc/doc/invoke.texi
@@ -2795,8 +2795,12 @@ have support for @option{-pthread}.
@item -fopenacc-dim=@var{geom}
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.
@opindex fopenmp
@cindex OpenMP parallel
diff --git a/gcc/omp-offload.cc b/gcc/omp-offload.cc
index 738b64f8c16..b2ba9042758 100644
--- a/gcc/omp-offload.cc
+++ b/gcc/omp-offload.cc
@@ -875,8 +875,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. */
@@ -911,14 +912,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 21eab80c776..bb18fa23628 100644
--- a/gcc/testsuite/ChangeLog.omp
+++ b/gcc/testsuite/ChangeLog.omp
@@ -1,3 +1,13 @@
+2018-10-05 Nathan Sidwell <nat...@acm.org>
+ Tom de Vries <tdevr...@suse.de>
+ Thomas Schwinge <tho...@codesourcery.com>
+ Julian Brown <jul...@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.
+
2018-10-22 Cesar Philippidis <ce...@codesourcery.com>
* g++.dg/goacc/loop-1.c: New test.
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 <nat...@acm.org>
+ Tom de Vries <tdevr...@suse.de>
+ Thomas Schwinge <tho...@codesourcery.com>
+ Julian Brown <jul...@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 <jnor...@codesourcery.com>
Cesar Philippidis <ce...@codesourcery.com>
Tom de Vries <t...@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;
+}
