diff options
| author | Tom de Vries <tdevries@suse.de> | 2019-01-12 22:17:42 +0000 |
|---|---|---|
| committer | Tom de Vries <vries@gcc.gnu.org> | 2019-01-12 22:17:42 +0000 |
| commit | 2b9d9e393766d2fa6e2dd5f361d0db14872cf261 (patch) | |
| tree | f7686faf57ee3b6e79c8818d2179d328251d77d4 | |
| parent | 6f7814d08fd89dc9f41807568ef26c3ffe498dba (diff) | |
| download | gcc-2b9d9e393766d2fa6e2dd5f361d0db14872cf261.zip gcc-2b9d9e393766d2fa6e2dd5f361d0db14872cf261.tar.gz gcc-2b9d9e393766d2fa6e2dd5f361d0db14872cf261.tar.bz2 | |
[nvptx] Enable large vectors
Allow vector_length clauses to accept values larger than warp size. Note that
this does not enable setting vector_length to values larger than warp size using
-fopenacc-dim.
2019-01-12 Tom de Vries <tdevries@suse.de>
* config/nvptx/nvptx.c (nvptx_goacc_validate_dims): Take larger vector
lengths into account.
* testsuite/libgomp.oacc-c-c++-common/vector-length-128-1.c: Expect
vector length to be 128.
* testsuite/libgomp.oacc-c-c++-common/parallel-dims.c: Expect vector
length 2097152 to be reduced to 1024 instead of 32.
From-SVN: r267889
| -rw-r--r-- | gcc/ChangeLog | 5 | ||||
| -rw-r--r-- | gcc/config/nvptx/nvptx.c | 2 | ||||
| -rw-r--r-- | libgomp/ChangeLog | 7 | ||||
| -rw-r--r-- | libgomp/testsuite/libgomp.oacc-c-c++-common/parallel-dims.c | 4 | ||||
| -rw-r--r-- | libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-1.c | 5 |
5 files changed, 17 insertions, 6 deletions
diff --git a/gcc/ChangeLog b/gcc/ChangeLog index a2735a5..6f18a34 100644 --- a/gcc/ChangeLog +++ b/gcc/ChangeLog @@ -1,3 +1,8 @@ +2019-01-12 Tom de Vries <tdevries@suse.de> + + * config/nvptx/nvptx.c (nvptx_goacc_validate_dims): Take larger vector + lengths into account. + 2019-01-12 Svante Signell <svante.signell@gmail.com> * config/i386/gnu.h (TARGET_THREAD_SSP_OFFSET): Define. diff --git a/gcc/config/nvptx/nvptx.c b/gcc/config/nvptx/nvptx.c index 1d97045..8d2740c 100644 --- a/gcc/config/nvptx/nvptx.c +++ b/gcc/config/nvptx/nvptx.c @@ -96,7 +96,7 @@ #define PTX_NUM_PER_WORKER_BARRIERS (PTX_CTA_NUM_BARRIERS - PTX_NUM_PER_CTA_BARRIERS) #define PTX_DEFAULT_VECTOR_LENGTH PTX_WARP_SIZE -#define PTX_MAX_VECTOR_LENGTH PTX_WARP_SIZE +#define PTX_MAX_VECTOR_LENGTH PTX_CTA_SIZE #define PTX_WORKER_LENGTH 32 #define PTX_DEFAULT_RUNTIME_DIM 0 /* Defer to runtime. */ diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog index a52fc92..dce5dce 100644 --- a/libgomp/ChangeLog +++ b/libgomp/ChangeLog @@ -1,3 +1,10 @@ +2019-01-12 Tom de Vries <tdevries@suse.de> + + * testsuite/libgomp.oacc-c-c++-common/vector-length-128-1.c: Expect + vector length to be 128. + * testsuite/libgomp.oacc-c-c++-common/parallel-dims.c: Expect vector + length 2097152 to be reduced to 1024 instead of 32. + 2019-01-11 Thomas Schwinge <thomas@codesourcery.com> James Norris <jnorris@codesourcery.com> diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/parallel-dims.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/parallel-dims.c index 4a98546..d7cd046 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/parallel-dims.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/parallel-dims.c @@ -350,7 +350,7 @@ int main () int gangs_min, gangs_max, workers_min, workers_max, vectors_min, vectors_max; gangs_min = workers_min = vectors_min = INT_MAX; gangs_max = workers_max = vectors_max = INT_MIN; -#pragma acc parallel copy (vectors_actual) /* { dg-warning "using vector_length \\(32\\), ignoring 2097152" "" { target openacc_nvidia_accel_configured } } */ \ +#pragma acc parallel copy (vectors_actual) /* { dg-warning "using vector_length \\(1024\\), ignoring 2097152" "" { target openacc_nvidia_accel_configured } } */ \ vector_length (VECTORS) { if (acc_on_device (acc_device_host)) @@ -361,7 +361,7 @@ int main () else if (acc_on_device (acc_device_nvidia)) { /* The GCC nvptx back end enforces vector_length (32). */ - vectors_actual = 32; + vectors_actual = 1024; } else __builtin_abort (); diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-1.c index fab5b0d..18d77cc 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-1.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-1.c @@ -33,7 +33,6 @@ main (void) return 0; } -/* { dg-prune-output "using vector_length \\(32\\), ignoring 128" } */ -/* { dg-final { scan-offload-tree-dump "__attribute__\\(\\(oacc function \\(1, 1, 32\\)" "oaccdevlow" } } */ -/* { dg-output "nvptx_exec: kernel main\\\$_omp_fn\\\$0: launch gangs=1, workers=1, vectors=32" } */ +/* { dg-final { scan-offload-tree-dump "__attribute__\\(\\(oacc function \\(1, 1, 128\\)" "oaccdevlow" } } */ +/* { dg-output "nvptx_exec: kernel main\\\$_omp_fn\\\$0: launch gangs=1, workers=1, vectors=128" } */ |