diff options
| author | Tom de Vries <tom@codesourcery.com> | 2018-04-05 08:36:37 +0000 |
|---|---|---|
| committer | Tom de Vries <vries@gcc.gnu.org> | 2018-04-05 08:36:37 +0000 |
| commit | 2ba16fd2eb40c96e41de967ca32e4dea4b5e45a1 (patch) | |
| tree | 26dcac04b2435d5bda5b7eba0accc7e71cccdde7 | |
| parent | 44780b91eb61d1a4ae3a6f94300d1b97a1d06557 (diff) | |
| download | gcc-2ba16fd2eb40c96e41de967ca32e4dea4b5e45a1.zip gcc-2ba16fd2eb40c96e41de967ca32e4dea4b5e45a1.tar.gz gcc-2ba16fd2eb40c96e41de967ca32e4dea4b5e45a1.tar.bz2 | |
[nvptx] Fix neutering of bb with only cond jump
2018-04-05 Tom de Vries <tom@codesourcery.com>
PR target/85204
* config/nvptx/nvptx.c (nvptx_single): Fix neutering of bb with only
cond jump.
* testsuite/libgomp.oacc-c-c++-common/broadcast-1.c: New test.
From-SVN: r259125
| -rw-r--r-- | gcc/ChangeLog | 6 | ||||
| -rw-r--r-- | gcc/config/nvptx/nvptx.c | 6 | ||||
| -rw-r--r-- | libgomp/ChangeLog | 5 | ||||
| -rw-r--r-- | libgomp/testsuite/libgomp.oacc-c-c++-common/broadcast-1.c | 49 |
4 files changed, 65 insertions, 1 deletions
diff --git a/gcc/ChangeLog b/gcc/ChangeLog index 6e299cb..517ac4e5 100644 --- a/gcc/ChangeLog +++ b/gcc/ChangeLog @@ -1,3 +1,9 @@ +2018-04-05 Tom de Vries <tom@codesourcery.com> + + PR target/85204 + * config/nvptx/nvptx.c (nvptx_single): Fix neutering of bb with only + cond jump. + 2018-04-05 Shiva Chen <shiva0217@gmail.com> Kito Cheng <kito.cheng@gmail.com> diff --git a/gcc/config/nvptx/nvptx.c b/gcc/config/nvptx/nvptx.c index b2b150f..a9a3053 100644 --- a/gcc/config/nvptx/nvptx.c +++ b/gcc/config/nvptx/nvptx.c @@ -4048,6 +4048,7 @@ nvptx_single (unsigned mask, basic_block from, basic_block to) /* Insert the vector test inside the worker test. */ unsigned mode; rtx_insn *before = tail; + rtx_insn *neuter_start = NULL; for (mode = GOMP_DIM_WORKER; mode <= GOMP_DIM_VECTOR; mode++) if (GOMP_DIM_MASK (mode) & skip_mask) { @@ -4065,7 +4066,10 @@ nvptx_single (unsigned mask, basic_block from, basic_block to) br = gen_br_true (pred, label); else br = gen_br_true_uni (pred, label); - emit_insn_before (br, head); + if (neuter_start) + neuter_start = emit_insn_after (br, neuter_start); + else + neuter_start = emit_insn_before (br, head); LABEL_NUSES (label)++; if (tail_branch) diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog index a5a5e06..ea28859 100644 --- a/libgomp/ChangeLog +++ b/libgomp/ChangeLog @@ -1,3 +1,8 @@ +2018-04-05 Tom de Vries <tom@codesourcery.com> + + PR target/85204 + * testsuite/libgomp.oacc-c-c++-common/broadcast-1.c: New test. + 2018-03-26 Tom de Vries <tom@codesourcery.com> PR tree-optimization/85063 diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/broadcast-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/broadcast-1.c new file mode 100644 index 0000000..ca0d37b --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/broadcast-1.c @@ -0,0 +1,49 @@ +/* Ensure that worker-vector state conditional expressions are + properly handled by the nvptx backend. */ + +#include <assert.h> +#include <math.h> + + +#define N 1024 + +int A[N][N] ; + +void test(int x) +{ +#pragma acc parallel num_gangs(16) num_workers(4) vector_length(32) copyout(A) + { +#pragma acc loop gang + for(int j=0;j<N;j++) + { + if (x==1) + { +#pragma acc loop worker vector + for(int i=0;i<N;i++) + A[i][j] = 1; + } + else + { +#pragma acc loop worker vector + for(int i=0;i<N;i++) + A[i][j] = -1; + } + } + } +} + + +int main(void) +{ + test (0); + for (int i = 0; i < N; i++) + for (int j = 0; j < N; j++) + assert (A[i][j] == -1); + + test (1); + for (int i = 0; i < N; i++) + for (int j = 0; j < N; j++) + assert (A[i][j] == 1); + + return 0; +} |