From 2b9d9e393766d2fa6e2dd5f361d0db14872cf261 Mon Sep 17 00:00:00 2001 From: Tom de Vries Date: Sat, 12 Jan 2019 22:17:42 +0000 Subject: [PATCH] [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 * 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 --- gcc/ChangeLog | 5 +++++ gcc/config/nvptx/nvptx.c | 2 +- libgomp/ChangeLog | 7 +++++++ .../testsuite/libgomp.oacc-c-c++-common/parallel-dims.c | 4 ++-- .../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 a2735a5e4c9..6f18a347b34 100644 --- a/gcc/ChangeLog +++ b/gcc/ChangeLog @@ -1,3 +1,8 @@ +2019-01-12 Tom de Vries + + * config/nvptx/nvptx.c (nvptx_goacc_validate_dims): Take larger vector + lengths into account. + 2019-01-12 Svante Signell * config/i386/gnu.h (TARGET_THREAD_SSP_OFFSET): Define. diff --git a/gcc/config/nvptx/nvptx.c b/gcc/config/nvptx/nvptx.c index 1d9704543d9..8d2740cd50f 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 a52fc92e118..dce5dcedf57 100644 --- a/libgomp/ChangeLog +++ b/libgomp/ChangeLog @@ -1,3 +1,10 @@ +2019-01-12 Tom de Vries + + * 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 James Norris 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 4a9854662cc..d7cd0461b53 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 fab5b0d25d1..18d77cc5ecb 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" } */ -- 2.30.2