From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: by sourceware.org (Postfix, from userid 1729) id EC59438485B5; Wed, 29 Jun 2022 14:34:30 +0000 (GMT) DKIM-Filter: OpenDKIM Filter v2.11.0 sourceware.org EC59438485B5 Content-Type: text/plain; charset="us-ascii" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit From: Kwok Yeung To: gcc-cvs@gcc.gnu.org Subject: [gcc/devel/omp/gcc-12] Default compute dimensions (compile time) X-Act-Checkin: gcc X-Git-Author: Julian Brown X-Git-Refname: refs/heads/devel/omp/gcc-12 X-Git-Oldrev: 30025fc576f5bca8901ab4cceb328e03f1f783ae X-Git-Newrev: d9075a4c4782a7d95acc6ebee9aad6ce1eabbccb Message-Id: <20220629143430.EC59438485B5@sourceware.org> Date: Wed, 29 Jun 2022 14:34:30 +0000 (GMT) X-BeenThere: gcc-cvs@gcc.gnu.org X-Mailman-Version: 2.1.29 Precedence: list List-Id: Gcc-cvs mailing list List-Unsubscribe: , List-Archive: List-Help: List-Subscribe: , X-List-Received-Date: Wed, 29 Jun 2022 14:34:31 -0000 https://gcc.gnu.org/g:d9075a4c4782a7d95acc6ebee9aad6ce1eabbccb commit d9075a4c4782a7d95acc6ebee9aad6ce1eabbccb Author: Julian Brown 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 Tom de Vries Thomas Schwinge Julian Brown 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 + Tom de Vries + Thomas Schwinge + Julian Brown + + * doc/invoke.texi (fopenacc-dim): Update. + * omp-offload.cc (oacc_parse_default_dims): Update. + 2019-09-20 Chung-Lin Tang Cesar Philippidis 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 + Tom de Vries + Thomas Schwinge + Julian Brown + + * 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 Cesar Philippidis 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 + Tom de Vries + Thomas Schwinge + Julian Brown + + * 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 Cesar Philippidis Tom de Vries 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; +}