From: Julian Brown Date: Tue, 26 Feb 2019 22:12:06 +0000 (-0800) Subject: Default compute dimensions (compile time) X-Git-Url: http://git.ipfire.org/cgi-bin/gitweb.cgi?a=commitdiff_plain;h=d9075a4c4782a7d95acc6ebee9aad6ce1eabbccb;p=thirdparty%2Fgcc.git 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 --git a/gcc/ChangeLog.omp b/gcc/ChangeLog.omp index ec14b9c60a00..ea0d696347ff 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 07b440190c3a..3b19eb80212d 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 78c2982da5e2..6131479b16f0 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 52910f0e4225..4108c8dedc62 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 9cf119bf89c7..bc2e0fd19b92 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 2a8d35d493de..6b1e7b224510 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 53b1fbe5039f..85fd50fb3344 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 42bcb0e8d63b..4249b404962c 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 753451587362..3179b778c690 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 000000000000..6c479e4eb258 --- /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 000000000000..20a022f27585 --- /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; +}